Cycles: Split device_opencl.cpp into multiple files for easier maintenance

There are no user-visible changes, just some internal restructuring.

Differential Revision: https://developer.blender.org/D2231
This commit is contained in:
Lukas Stockner 2016-09-14 23:47:54 +02:00
parent 1fad269d07
commit 9ea71bc674
Notes: blender-bot 2023-02-14 11:21:40 +01:00
Referenced by issue #50135, Particle - Hair Objects - Hair objects are generated at origin point, not on mesh
Referenced by issue #50104, Cycles - Material compilation randomly fails
Referenced by issue #49630, regression MIS Split kernel OpenCL
Referenced by issue #49623, Immediately crash trying to render attached file in Cycles
Referenced by issue #49592, Crash on "Load Factory Setting"
7 changed files with 3364 additions and 3258 deletions

View File

@ -36,6 +36,15 @@ set(SRC
device_task.cpp
)
set(SRC_OPENCL
opencl/opencl.h
opencl/opencl_base.cpp
opencl/opencl_mega.cpp
opencl/opencl_split.cpp
opencl/opencl_util.cpp
)
if(WITH_CYCLES_NETWORK)
list(APPEND SRC
device_network.cpp
@ -67,4 +76,4 @@ endif()
include_directories(${INC})
include_directories(SYSTEM ${INC_SYS})
add_library(cycles_device ${SRC} ${SRC_HEADERS})
add_library(cycles_device ${SRC} ${SRC_OPENCL} ${SRC_HEADERS})

File diff suppressed because it is too large Load Diff

View File

@ -0,0 +1,397 @@
/*
* Copyright 2011-2013 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifdef WITH_OPENCL
#include "clew.h"
#include "device.h"
#include "util_map.h"
#include "util_param.h"
#include "util_string.h"
CCL_NAMESPACE_BEGIN
#define CL_MEM_PTR(p) ((cl_mem)(uintptr_t)(p))
/* Macro declarations used with split kernel */
/* Macro to enable/disable work-stealing */
#define __WORK_STEALING__
#define SPLIT_KERNEL_LOCAL_SIZE_X 64
#define SPLIT_KERNEL_LOCAL_SIZE_Y 1
/* This value may be tuned according to the scene we are rendering.
*
* Modifying PATH_ITER_INC_FACTOR value proportional to number of expected
* ray-bounces will improve performance.
*/
#define PATH_ITER_INC_FACTOR 8
/* When allocate global memory in chunks. We may not be able to
* allocate exactly "CL_DEVICE_MAX_MEM_ALLOC_SIZE" bytes in chunks;
* Since some bytes may be needed for aligning chunks of memory;
* This is the amount of memory that we dedicate for that purpose.
*/
#define DATA_ALLOCATION_MEM_FACTOR 5000000 //5MB
struct OpenCLPlatformDevice {
OpenCLPlatformDevice(cl_platform_id platform_id,
const string& platform_name,
cl_device_id device_id,
cl_device_type device_type,
const string& device_name)
: platform_id(platform_id),
platform_name(platform_name),
device_id(device_id),
device_type(device_type),
device_name(device_name) {}
cl_platform_id platform_id;
string platform_name;
cl_device_id device_id;
cl_device_type device_type;
string device_name;
};
/* Contains all static OpenCL helper functions. */
class OpenCLInfo
{
public:
static cl_device_type device_type();
static bool use_debug();
static bool kernel_use_advanced_shading(const string& platform_name);
static bool kernel_use_split(const string& platform_name,
const cl_device_type device_type);
static bool device_supported(const string& platform_name,
const cl_device_id device_id);
static bool platform_version_check(cl_platform_id platform,
string *error = NULL);
static bool device_version_check(cl_device_id device,
string *error = NULL);
static void get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices,
bool force_all = false);
};
/* Thread safe cache for contexts and programs.
*/
class OpenCLCache
{
struct Slot
{
struct ProgramEntry
{
ProgramEntry();
ProgramEntry(const ProgramEntry& rhs);
~ProgramEntry();
cl_program program;
thread_mutex *mutex;
};
Slot();
Slot(const Slot& rhs);
~Slot();
thread_mutex *context_mutex;
cl_context context;
typedef map<ustring, ProgramEntry> EntryMap;
EntryMap programs;
};
/* key is combination of platform ID and device ID */
typedef pair<cl_platform_id, cl_device_id> PlatformDevicePair;
/* map of Slot objects */
typedef map<PlatformDevicePair, Slot> CacheMap;
CacheMap cache;
/* MD5 hash of the kernel source. */
string kernel_md5;
thread_mutex cache_lock;
thread_mutex kernel_md5_lock;
/* lazy instantiate */
static OpenCLCache& global_instance();
public:
enum ProgramName {
OCL_DEV_BASE_PROGRAM,
OCL_DEV_MEGAKERNEL_PROGRAM,
};
/* Lookup context in the cache. If this returns NULL, slot_locker
* will be holding a lock for the cache. slot_locker should refer to a
* default constructed thread_scoped_lock. */
static cl_context get_context(cl_platform_id platform,
cl_device_id device,
thread_scoped_lock& slot_locker);
/* Same as above. */
static cl_program get_program(cl_platform_id platform,
cl_device_id device,
ustring key,
thread_scoped_lock& slot_locker);
/* Store context in the cache. You MUST have tried to get the item before storing to it. */
static void store_context(cl_platform_id platform,
cl_device_id device,
cl_context context,
thread_scoped_lock& slot_locker);
/* Same as above. */
static void store_program(cl_platform_id platform,
cl_device_id device,
cl_program program,
ustring key,
thread_scoped_lock& slot_locker);
static string get_kernel_md5();
};
#define opencl_assert(stmt) \
{ \
cl_int err = stmt; \
\
if(err != CL_SUCCESS) { \
string message = string_printf("OpenCL error: %s in %s", clewErrorString(err), #stmt); \
if(error_msg == "") \
error_msg = message; \
fprintf(stderr, "%s\n", message.c_str()); \
} \
} (void)0
class OpenCLDeviceBase : public Device
{
public:
DedicatedTaskPool task_pool;
cl_context cxContext;
cl_command_queue cqCommandQueue;
cl_platform_id cpPlatform;
cl_device_id cdDevice;
cl_int ciErr;
class OpenCLProgram {
public:
OpenCLProgram() : loaded(false), device(NULL) {}
OpenCLProgram(OpenCLDeviceBase *device,
string program_name,
string kernel_name,
string kernel_build_options);
~OpenCLProgram();
void add_kernel(ustring name);
void load();
bool is_loaded() { return loaded; }
string get_log() { return log; }
void report_error();
cl_kernel operator()();
cl_kernel operator()(ustring name);
void release();
private:
bool build_kernel(const string *debug_src);
bool compile_kernel(const string *debug_src);
bool load_binary(const string& clbin, const string *debug_src = NULL);
bool save_binary(const string& clbin);
bool loaded;
cl_program program;
OpenCLDeviceBase *device;
/* Used for the OpenCLCache key. */
string program_name;
string kernel_file, kernel_build_options, device_md5;
string error_msg, output_msg;
string log;
map<ustring, cl_kernel> kernels;
};
OpenCLProgram base_program;
typedef map<string, device_vector<uchar>*> ConstMemMap;
typedef map<string, device_ptr> MemMap;
ConstMemMap const_mem_map;
MemMap mem_map;
device_ptr null_mem;
bool device_initialized;
string platform_name;
bool opencl_error(cl_int err);
void opencl_error(const string& message);
void opencl_assert_err(cl_int err, const char* where);
OpenCLDeviceBase(DeviceInfo& info, Stats &stats, bool background_);
~OpenCLDeviceBase();
static void CL_CALLBACK context_notify_callback(const char *err_info,
const void * /*private_info*/, size_t /*cb*/, void *user_data);
bool opencl_version_check();
string device_md5_hash(string kernel_custom_build_options = "");
bool load_kernels(const DeviceRequestedFeatures& requested_features);
/* Has to be implemented by the real device classes.
* The base device will then load all these programs. */
virtual void load_kernels(const DeviceRequestedFeatures& requested_features,
vector<OpenCLProgram*> &programs) = 0;
void mem_alloc(device_memory& mem, MemoryType type);
void mem_copy_to(device_memory& mem);
void mem_copy_from(device_memory& mem, int y, int w, int h, int elem);
void mem_zero(device_memory& mem);
void mem_free(device_memory& mem);
void const_copy_to(const char *name, void *host, size_t size);
void tex_alloc(const char *name,
device_memory& mem,
InterpolationType /*interpolation*/,
ExtensionType /*extension*/);
void tex_free(device_memory& mem);
size_t global_size_round_up(int group_size, int global_size);
void enqueue_kernel(cl_kernel kernel, size_t w, size_t h);
void set_kernel_arg_mem(cl_kernel kernel, cl_uint *narg, const char *name);
void film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half);
void shader(DeviceTask& task);
class OpenCLDeviceTask : public DeviceTask {
public:
OpenCLDeviceTask(OpenCLDeviceBase *device, DeviceTask& task)
: DeviceTask(task)
{
run = function_bind(&OpenCLDeviceBase::thread_run,
device,
this);
}
};
int get_split_task_count(DeviceTask& /*task*/)
{
return 1;
}
void task_add(DeviceTask& task)
{
task_pool.push(new OpenCLDeviceTask(this, task));
}
void task_wait()
{
task_pool.wait();
}
void task_cancel()
{
task_pool.cancel();
}
virtual void thread_run(DeviceTask * /*task*/) = 0;
protected:
string kernel_build_options(const string *debug_src = NULL);
class ArgumentWrapper {
public:
ArgumentWrapper() : size(0), pointer(NULL) {}
template <typename T>
ArgumentWrapper(T& argument) : size(sizeof(argument)),
pointer(&argument) { }
ArgumentWrapper(int argument) : size(sizeof(int)),
int_value(argument),
pointer(&int_value) { }
ArgumentWrapper(float argument) : size(sizeof(float)),
float_value(argument),
pointer(&float_value) { }
size_t size;
int int_value;
float float_value;
void *pointer;
};
/* TODO(sergey): In the future we can use variadic templates, once
* C++0x is allowed. Should allow to clean this up a bit.
*/
int kernel_set_args(cl_kernel kernel,
int start_argument_index,
const ArgumentWrapper& arg1 = ArgumentWrapper(),
const ArgumentWrapper& arg2 = ArgumentWrapper(),
const ArgumentWrapper& arg3 = ArgumentWrapper(),
const ArgumentWrapper& arg4 = ArgumentWrapper(),
const ArgumentWrapper& arg5 = ArgumentWrapper(),
const ArgumentWrapper& arg6 = ArgumentWrapper(),
const ArgumentWrapper& arg7 = ArgumentWrapper(),
const ArgumentWrapper& arg8 = ArgumentWrapper(),
const ArgumentWrapper& arg9 = ArgumentWrapper(),
const ArgumentWrapper& arg10 = ArgumentWrapper(),
const ArgumentWrapper& arg11 = ArgumentWrapper(),
const ArgumentWrapper& arg12 = ArgumentWrapper(),
const ArgumentWrapper& arg13 = ArgumentWrapper(),
const ArgumentWrapper& arg14 = ArgumentWrapper(),
const ArgumentWrapper& arg15 = ArgumentWrapper(),
const ArgumentWrapper& arg16 = ArgumentWrapper(),
const ArgumentWrapper& arg17 = ArgumentWrapper(),
const ArgumentWrapper& arg18 = ArgumentWrapper(),
const ArgumentWrapper& arg19 = ArgumentWrapper(),
const ArgumentWrapper& arg20 = ArgumentWrapper(),
const ArgumentWrapper& arg21 = ArgumentWrapper(),
const ArgumentWrapper& arg22 = ArgumentWrapper(),
const ArgumentWrapper& arg23 = ArgumentWrapper(),
const ArgumentWrapper& arg24 = ArgumentWrapper(),
const ArgumentWrapper& arg25 = ArgumentWrapper(),
const ArgumentWrapper& arg26 = ArgumentWrapper(),
const ArgumentWrapper& arg27 = ArgumentWrapper(),
const ArgumentWrapper& arg28 = ArgumentWrapper(),
const ArgumentWrapper& arg29 = ArgumentWrapper(),
const ArgumentWrapper& arg30 = ArgumentWrapper(),
const ArgumentWrapper& arg31 = ArgumentWrapper(),
const ArgumentWrapper& arg32 = ArgumentWrapper(),
const ArgumentWrapper& arg33 = ArgumentWrapper());
void release_kernel_safe(cl_kernel kernel);
void release_mem_object_safe(cl_mem mem);
void release_program_safe(cl_program program);
/* ** Those guys are for workign around some compiler-specific bugs ** */
virtual cl_program load_cached_kernel(
ustring key,
thread_scoped_lock& cache_locker);
virtual void store_cached_kernel(
cl_program program,
ustring key,
thread_scoped_lock& cache_locker);
virtual string build_options_for_base_program(
const DeviceRequestedFeatures& /*requested_features*/);
};
Device *opencl_create_mega_device(DeviceInfo& info, Stats& stats, bool background);
Device *opencl_create_split_device(DeviceInfo& info, Stats& stats, bool background);
CCL_NAMESPACE_END
#endif

View File

@ -0,0 +1,728 @@
/*
* Copyright 2011-2013 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifdef WITH_OPENCL
#include "opencl.h"
#include "kernel_types.h"
#include "util_foreach.h"
#include "util_logging.h"
#include "util_md5.h"
#include "util_path.h"
#include "util_time.h"
CCL_NAMESPACE_BEGIN
bool OpenCLDeviceBase::opencl_error(cl_int err)
{
if(err != CL_SUCCESS) {
string message = string_printf("OpenCL error (%d): %s", err, clewErrorString(err));
if(error_msg == "")
error_msg = message;
fprintf(stderr, "%s\n", message.c_str());
return true;
}
return false;
}
void OpenCLDeviceBase::opencl_error(const string& message)
{
if(error_msg == "")
error_msg = message;
fprintf(stderr, "%s\n", message.c_str());
}
void OpenCLDeviceBase::opencl_assert_err(cl_int err, const char* where)
{
if(err != CL_SUCCESS) {
string message = string_printf("OpenCL error (%d): %s in %s", err, clewErrorString(err), where);
if(error_msg == "")
error_msg = message;
fprintf(stderr, "%s\n", message.c_str());
#ifndef NDEBUG
abort();
#endif
}
}
OpenCLDeviceBase::OpenCLDeviceBase(DeviceInfo& info, Stats &stats, bool background_)
: Device(info, stats, background_)
{
cpPlatform = NULL;
cdDevice = NULL;
cxContext = NULL;
cqCommandQueue = NULL;
null_mem = 0;
device_initialized = false;
vector<OpenCLPlatformDevice> usable_devices;
OpenCLInfo::get_usable_devices(&usable_devices);
if(usable_devices.size() == 0) {
opencl_error("OpenCL: no devices found.");
return;
}
assert(info.num < usable_devices.size());
OpenCLPlatformDevice& platform_device = usable_devices[info.num];
cpPlatform = platform_device.platform_id;
cdDevice = platform_device.device_id;
platform_name = platform_device.platform_name;
VLOG(2) << "Creating new Cycles device for OpenCL platform "
<< platform_name << ", device "
<< platform_device.device_name << ".";
{
/* try to use cached context */
thread_scoped_lock cache_locker;
cxContext = OpenCLCache::get_context(cpPlatform, cdDevice, cache_locker);
if(cxContext == NULL) {
/* create context properties array to specify platform */
const cl_context_properties context_props[] = {
CL_CONTEXT_PLATFORM, (cl_context_properties)cpPlatform,
0, 0
};
/* create context */
cxContext = clCreateContext(context_props, 1, &cdDevice,
context_notify_callback, cdDevice, &ciErr);
if(opencl_error(ciErr)) {
opencl_error("OpenCL: clCreateContext failed");
return;
}
/* cache it */
OpenCLCache::store_context(cpPlatform, cdDevice, cxContext, cache_locker);
}
}
cqCommandQueue = clCreateCommandQueue(cxContext, cdDevice, 0, &ciErr);
if(opencl_error(ciErr))
return;
null_mem = (device_ptr)clCreateBuffer(cxContext, CL_MEM_READ_ONLY, 1, NULL, &ciErr);
if(opencl_error(ciErr))
return;
fprintf(stderr, "Device init success\n");
device_initialized = true;
}
OpenCLDeviceBase::~OpenCLDeviceBase()
{
task_pool.stop();
if(null_mem)
clReleaseMemObject(CL_MEM_PTR(null_mem));
ConstMemMap::iterator mt;
for(mt = const_mem_map.begin(); mt != const_mem_map.end(); mt++) {
mem_free(*(mt->second));
delete mt->second;
}
base_program.release();
if(cqCommandQueue)
clReleaseCommandQueue(cqCommandQueue);
if(cxContext)
clReleaseContext(cxContext);
}
void CL_CALLBACK OpenCLDeviceBase::context_notify_callback(const char *err_info,
const void * /*private_info*/, size_t /*cb*/, void *user_data)
{
char name[256];
clGetDeviceInfo((cl_device_id)user_data, CL_DEVICE_NAME, sizeof(name), &name, NULL);
fprintf(stderr, "OpenCL error (%s): %s\n", name, err_info);
}
bool OpenCLDeviceBase::opencl_version_check()
{
string error;
if(!OpenCLInfo::platform_version_check(cpPlatform, &error)) {
opencl_error(error);
return false;
}
if(!OpenCLInfo::device_version_check(cdDevice, &error)) {
opencl_error(error);
return false;
}
return true;
}
string OpenCLDeviceBase::device_md5_hash(string kernel_custom_build_options)
{
MD5Hash md5;
char version[256], driver[256], name[256], vendor[256];
clGetPlatformInfo(cpPlatform, CL_PLATFORM_VENDOR, sizeof(vendor), &vendor, NULL);
clGetDeviceInfo(cdDevice, CL_DEVICE_VERSION, sizeof(version), &version, NULL);
clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(name), &name, NULL);
clGetDeviceInfo(cdDevice, CL_DRIVER_VERSION, sizeof(driver), &driver, NULL);
md5.append((uint8_t*)vendor, strlen(vendor));
md5.append((uint8_t*)version, strlen(version));
md5.append((uint8_t*)name, strlen(name));
md5.append((uint8_t*)driver, strlen(driver));
string options = kernel_build_options();
options += kernel_custom_build_options;
md5.append((uint8_t*)options.c_str(), options.size());
return md5.get_hex();
}
bool OpenCLDeviceBase::load_kernels(const DeviceRequestedFeatures& requested_features)
{
/* Verify if device was initialized. */
if(!device_initialized) {
fprintf(stderr, "OpenCL: failed to initialize device.\n");
return false;
}
/* Verify we have right opencl version. */
if(!opencl_version_check())
return false;
base_program = OpenCLProgram(this, "base", "kernel.cl", build_options_for_base_program(requested_features));
base_program.add_kernel(ustring("convert_to_byte"));
base_program.add_kernel(ustring("convert_to_half_float"));
base_program.add_kernel(ustring("shader"));
base_program.add_kernel(ustring("bake"));
vector<OpenCLProgram*> programs;
programs.push_back(&base_program);
/* Call actual class to fill the vector with its programs. */
load_kernels(requested_features, programs);
TaskPool task_pool;
foreach(OpenCLProgram *program, programs) {
task_pool.push(function_bind(&OpenCLProgram::load, program));
}
task_pool.wait_work();
foreach(OpenCLProgram *program, programs) {
VLOG(2) << program->get_log();
if(!program->is_loaded()) {
program->report_error();
return false;
}
}
return true;
}
void OpenCLDeviceBase::mem_alloc(device_memory& mem, MemoryType type)
{
size_t size = mem.memory_size();
cl_mem_flags mem_flag;
void *mem_ptr = NULL;
if(type == MEM_READ_ONLY)
mem_flag = CL_MEM_READ_ONLY;
else if(type == MEM_WRITE_ONLY)
mem_flag = CL_MEM_WRITE_ONLY;
else
mem_flag = CL_MEM_READ_WRITE;
/* Zero-size allocation might be invoked by render, but not really
* supported by OpenCL. Using NULL as device pointer also doesn't really
* work for some reason, so for the time being we'll use special case
* will null_mem buffer.
*/
if(size != 0) {
mem.device_pointer = (device_ptr)clCreateBuffer(cxContext,
mem_flag,
size,
mem_ptr,
&ciErr);
opencl_assert_err(ciErr, "clCreateBuffer");
}
else {
mem.device_pointer = null_mem;
}
stats.mem_alloc(size);
mem.device_size = size;
}
void OpenCLDeviceBase::mem_copy_to(device_memory& mem)
{
/* this is blocking */
size_t size = mem.memory_size();
if(size != 0) {
opencl_assert(clEnqueueWriteBuffer(cqCommandQueue,
CL_MEM_PTR(mem.device_pointer),
CL_TRUE,
0,
size,
(void*)mem.data_pointer,
0,
NULL, NULL));
}
}
void OpenCLDeviceBase::mem_copy_from(device_memory& mem, int y, int w, int h, int elem)
{
size_t offset = elem*y*w;
size_t size = elem*w*h;
assert(size != 0);
opencl_assert(clEnqueueReadBuffer(cqCommandQueue,
CL_MEM_PTR(mem.device_pointer),
CL_TRUE,
offset,
size,
(uchar*)mem.data_pointer + offset,
0,
NULL, NULL));
}
void OpenCLDeviceBase::mem_zero(device_memory& mem)
{
if(mem.device_pointer) {
memset((void*)mem.data_pointer, 0, mem.memory_size());
mem_copy_to(mem);
}
}
void OpenCLDeviceBase::mem_free(device_memory& mem)
{
if(mem.device_pointer) {
if(mem.device_pointer != null_mem) {
opencl_assert(clReleaseMemObject(CL_MEM_PTR(mem.device_pointer)));
}
mem.device_pointer = 0;
stats.mem_free(mem.device_size);
mem.device_size = 0;
}
}
void OpenCLDeviceBase::const_copy_to(const char *name, void *host, size_t size)
{
ConstMemMap::iterator i = const_mem_map.find(name);
if(i == const_mem_map.end()) {
device_vector<uchar> *data = new device_vector<uchar>();
data->copy((uchar*)host, size);
mem_alloc(*data, MEM_READ_ONLY);
i = const_mem_map.insert(ConstMemMap::value_type(name, data)).first;
}
else {
device_vector<uchar> *data = i->second;
data->copy((uchar*)host, size);
}
mem_copy_to(*i->second);
}
void OpenCLDeviceBase::tex_alloc(const char *name,
device_memory& mem,
InterpolationType /*interpolation*/,
ExtensionType /*extension*/)
{
VLOG(1) << "Texture allocate: " << name << ", "
<< string_human_readable_number(mem.memory_size()) << " bytes. ("
<< string_human_readable_size(mem.memory_size()) << ")";
mem_alloc(mem, MEM_READ_ONLY);
mem_copy_to(mem);
assert(mem_map.find(name) == mem_map.end());
mem_map.insert(MemMap::value_type(name, mem.device_pointer));
}
void OpenCLDeviceBase::tex_free(device_memory& mem)
{
if(mem.device_pointer) {
foreach(const MemMap::value_type& value, mem_map) {
if(value.second == mem.device_pointer) {
mem_map.erase(value.first);
break;
}
}
mem_free(mem);
}
}
size_t OpenCLDeviceBase::global_size_round_up(int group_size, int global_size)
{
int r = global_size % group_size;
return global_size + ((r == 0)? 0: group_size - r);
}
void OpenCLDeviceBase::enqueue_kernel(cl_kernel kernel, size_t w, size_t h)
{
size_t workgroup_size, max_work_items[3];
clGetKernelWorkGroupInfo(kernel, cdDevice,
CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL);
clGetDeviceInfo(cdDevice,
CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(size_t)*3, max_work_items, NULL);
/* Try to divide evenly over 2 dimensions. */
size_t sqrt_workgroup_size = max((size_t)sqrt((double)workgroup_size), 1);
size_t local_size[2] = {sqrt_workgroup_size, sqrt_workgroup_size};
/* Some implementations have max size 1 on 2nd dimension. */
if(local_size[1] > max_work_items[1]) {
local_size[0] = workgroup_size/max_work_items[1];
local_size[1] = max_work_items[1];
}
size_t global_size[2] = {global_size_round_up(local_size[0], w),
global_size_round_up(local_size[1], h)};
/* Vertical size of 1 is coming from bake/shade kernels where we should
* not round anything up because otherwise we'll either be doing too
* much work per pixel (if we don't check global ID on Y axis) or will
* be checking for global ID to always have Y of 0.
*/
if (h == 1) {
global_size[h] = 1;
}
/* run kernel */
opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, kernel, 2, NULL, global_size, NULL, 0, NULL, NULL));
opencl_assert(clFlush(cqCommandQueue));
}
void OpenCLDeviceBase::set_kernel_arg_mem(cl_kernel kernel, cl_uint *narg, const char *name)
{
cl_mem ptr;
MemMap::iterator i = mem_map.find(name);
if(i != mem_map.end()) {
ptr = CL_MEM_PTR(i->second);
}
else {
/* work around NULL not working, even though the spec says otherwise */
ptr = CL_MEM_PTR(null_mem);
}
opencl_assert(clSetKernelArg(kernel, (*narg)++, sizeof(ptr), (void*)&ptr));
}
void OpenCLDeviceBase::film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half)
{
/* cast arguments to cl types */
cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
cl_mem d_rgba = (rgba_byte)? CL_MEM_PTR(rgba_byte): CL_MEM_PTR(rgba_half);
cl_mem d_buffer = CL_MEM_PTR(buffer);
cl_int d_x = task.x;
cl_int d_y = task.y;
cl_int d_w = task.w;
cl_int d_h = task.h;
cl_float d_sample_scale = 1.0f/(task.sample + 1);
cl_int d_offset = task.offset;
cl_int d_stride = task.stride;
cl_kernel ckFilmConvertKernel = (rgba_byte)? base_program(ustring("convert_to_byte")): base_program(ustring("convert_to_half_float"));
cl_uint start_arg_index =
kernel_set_args(ckFilmConvertKernel,
0,
d_data,
d_rgba,
d_buffer);
#define KERNEL_TEX(type, ttype, name) \
set_kernel_arg_mem(ckFilmConvertKernel, &start_arg_index, #name);
#include "kernel_textures.h"
#undef KERNEL_TEX
start_arg_index += kernel_set_args(ckFilmConvertKernel,
start_arg_index,
d_sample_scale,
d_x,
d_y,
d_w,
d_h,
d_offset,
d_stride);
enqueue_kernel(ckFilmConvertKernel, d_w, d_h);
}
void OpenCLDeviceBase::shader(DeviceTask& task)
{
/* cast arguments to cl types */
cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
cl_mem d_input = CL_MEM_PTR(task.shader_input);
cl_mem d_output = CL_MEM_PTR(task.shader_output);
cl_mem d_output_luma = CL_MEM_PTR(task.shader_output_luma);
cl_int d_shader_eval_type = task.shader_eval_type;
cl_int d_shader_filter = task.shader_filter;
cl_int d_shader_x = task.shader_x;
cl_int d_shader_w = task.shader_w;
cl_int d_offset = task.offset;
cl_kernel kernel;
if(task.shader_eval_type >= SHADER_EVAL_BAKE)
kernel = base_program(ustring("shader"));
else
kernel = base_program(ustring("bake"));
cl_uint start_arg_index =
kernel_set_args(kernel,
0,
d_data,
d_input,
d_output);
if(task.shader_eval_type < SHADER_EVAL_BAKE) {
start_arg_index += kernel_set_args(kernel,
start_arg_index,
d_output_luma);
}
#define KERNEL_TEX(type, ttype, name) \
set_kernel_arg_mem(kernel, &start_arg_index, #name);
#include "kernel_textures.h"
#undef KERNEL_TEX
start_arg_index += kernel_set_args(kernel,
start_arg_index,
d_shader_eval_type);
if(task.shader_eval_type >= SHADER_EVAL_BAKE) {
start_arg_index += kernel_set_args(kernel,
start_arg_index,
d_shader_filter);
}
start_arg_index += kernel_set_args(kernel,
start_arg_index,
d_shader_x,
d_shader_w,
d_offset);
for(int sample = 0; sample < task.num_samples; sample++) {
if(task.get_cancel())
break;
kernel_set_args(kernel, start_arg_index, sample);
enqueue_kernel(kernel, task.shader_w, 1);
clFinish(cqCommandQueue);
task.update_progress(NULL);
}
}
string OpenCLDeviceBase::kernel_build_options(const string *debug_src)
{
string build_options = "-cl-fast-relaxed-math ";
if(platform_name == "NVIDIA CUDA") {
build_options += "-D__KERNEL_OPENCL_NVIDIA__ "
"-cl-nv-maxrregcount=32 "
"-cl-nv-verbose ";
uint compute_capability_major, compute_capability_minor;
clGetDeviceInfo(cdDevice, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV,
sizeof(cl_uint), &compute_capability_major, NULL);
clGetDeviceInfo(cdDevice, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV,
sizeof(cl_uint), &compute_capability_minor, NULL);
build_options += string_printf("-D__COMPUTE_CAPABILITY__=%u ",
compute_capability_major * 100 +
compute_capability_minor * 10);
}
else if(platform_name == "Apple")
build_options += "-D__KERNEL_OPENCL_APPLE__ ";
else if(platform_name == "AMD Accelerated Parallel Processing")
build_options += "-D__KERNEL_OPENCL_AMD__ ";
else if(platform_name == "Intel(R) OpenCL") {
build_options += "-D__KERNEL_OPENCL_INTEL_CPU__ ";
/* Options for gdb source level kernel debugging.
* this segfaults on linux currently.
*/
if(OpenCLInfo::use_debug() && debug_src)
build_options += "-g -s \"" + *debug_src + "\" ";
}
if(OpenCLInfo::use_debug())
build_options += "-D__KERNEL_OPENCL_DEBUG__ ";
#ifdef WITH_CYCLES_DEBUG
build_options += "-D__KERNEL_DEBUG__ ";
#endif
return build_options;
}
/* TODO(sergey): In the future we can use variadic templates, once
* C++0x is allowed. Should allow to clean this up a bit.
*/
int OpenCLDeviceBase::kernel_set_args(cl_kernel kernel,
int start_argument_index,
const ArgumentWrapper& arg1,
const ArgumentWrapper& arg2,
const ArgumentWrapper& arg3,
const ArgumentWrapper& arg4,
const ArgumentWrapper& arg5,
const ArgumentWrapper& arg6,
const ArgumentWrapper& arg7,
const ArgumentWrapper& arg8,
const ArgumentWrapper& arg9,
const ArgumentWrapper& arg10,
const ArgumentWrapper& arg11,
const ArgumentWrapper& arg12,
const ArgumentWrapper& arg13,
const ArgumentWrapper& arg14,
const ArgumentWrapper& arg15,
const ArgumentWrapper& arg16,
const ArgumentWrapper& arg17,
const ArgumentWrapper& arg18,
const ArgumentWrapper& arg19,
const ArgumentWrapper& arg20,
const ArgumentWrapper& arg21,
const ArgumentWrapper& arg22,
const ArgumentWrapper& arg23,
const ArgumentWrapper& arg24,
const ArgumentWrapper& arg25,
const ArgumentWrapper& arg26,
const ArgumentWrapper& arg27,
const ArgumentWrapper& arg28,
const ArgumentWrapper& arg29,
const ArgumentWrapper& arg30,
const ArgumentWrapper& arg31,
const ArgumentWrapper& arg32,
const ArgumentWrapper& arg33)
{
int current_arg_index = 0;
#define FAKE_VARARG_HANDLE_ARG(arg) \
do { \
if(arg.pointer != NULL) { \
opencl_assert(clSetKernelArg( \
kernel, \
start_argument_index + current_arg_index, \
arg.size, arg.pointer)); \
++current_arg_index; \
} \
else { \
return current_arg_index; \
} \
} while(false)
FAKE_VARARG_HANDLE_ARG(arg1);
FAKE_VARARG_HANDLE_ARG(arg2);
FAKE_VARARG_HANDLE_ARG(arg3);
FAKE_VARARG_HANDLE_ARG(arg4);
FAKE_VARARG_HANDLE_ARG(arg5);
FAKE_VARARG_HANDLE_ARG(arg6);
FAKE_VARARG_HANDLE_ARG(arg7);
FAKE_VARARG_HANDLE_ARG(arg8);
FAKE_VARARG_HANDLE_ARG(arg9);
FAKE_VARARG_HANDLE_ARG(arg10);
FAKE_VARARG_HANDLE_ARG(arg11);
FAKE_VARARG_HANDLE_ARG(arg12);
FAKE_VARARG_HANDLE_ARG(arg13);
FAKE_VARARG_HANDLE_ARG(arg14);
FAKE_VARARG_HANDLE_ARG(arg15);
FAKE_VARARG_HANDLE_ARG(arg16);
FAKE_VARARG_HANDLE_ARG(arg17);
FAKE_VARARG_HANDLE_ARG(arg18);
FAKE_VARARG_HANDLE_ARG(arg19);
FAKE_VARARG_HANDLE_ARG(arg20);
FAKE_VARARG_HANDLE_ARG(arg21);
FAKE_VARARG_HANDLE_ARG(arg22);
FAKE_VARARG_HANDLE_ARG(arg23);
FAKE_VARARG_HANDLE_ARG(arg24);
FAKE_VARARG_HANDLE_ARG(arg25);
FAKE_VARARG_HANDLE_ARG(arg26);
FAKE_VARARG_HANDLE_ARG(arg27);
FAKE_VARARG_HANDLE_ARG(arg28);
FAKE_VARARG_HANDLE_ARG(arg29);
FAKE_VARARG_HANDLE_ARG(arg30);
FAKE_VARARG_HANDLE_ARG(arg31);
FAKE_VARARG_HANDLE_ARG(arg32);
FAKE_VARARG_HANDLE_ARG(arg33);
#undef FAKE_VARARG_HANDLE_ARG
return current_arg_index;
}
void OpenCLDeviceBase::release_kernel_safe(cl_kernel kernel)
{
if(kernel) {
clReleaseKernel(kernel);
}
}
void OpenCLDeviceBase::release_mem_object_safe(cl_mem mem)
{
if(mem != NULL) {
clReleaseMemObject(mem);
}
}
void OpenCLDeviceBase::release_program_safe(cl_program program)
{
if(program) {
clReleaseProgram(program);
}
}
/* ** Those guys are for workign around some compiler-specific bugs ** */
cl_program OpenCLDeviceBase::load_cached_kernel(
ustring key,
thread_scoped_lock& cache_locker)
{
return OpenCLCache::get_program(cpPlatform,
cdDevice,
key,
cache_locker);
}
void OpenCLDeviceBase::store_cached_kernel(
cl_program program,
ustring key,
thread_scoped_lock& cache_locker)
{
OpenCLCache::store_program(cpPlatform,
cdDevice,
program,
key,
cache_locker);
}
string OpenCLDeviceBase::build_options_for_base_program(
const DeviceRequestedFeatures& /*requested_features*/)
{
/* TODO(sergey): By default we compile all features, meaning
* mega kernel is not getting feature-based optimizations.
*
* Ideally we need always compile kernel with as less features
* enabled as possible to keep performance at it's max.
*/
return "";
}
CCL_NAMESPACE_END
#endif

View File

@ -0,0 +1,150 @@
/*
* Copyright 2011-2013 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifdef WITH_OPENCL
#include "opencl.h"
#include "buffers.h"
#include "kernel_types.h"
#include "util_md5.h"
#include "util_path.h"
#include "util_time.h"
CCL_NAMESPACE_BEGIN
class OpenCLDeviceMegaKernel : public OpenCLDeviceBase
{
public:
OpenCLProgram path_trace_program;
OpenCLDeviceMegaKernel(DeviceInfo& info, Stats &stats, bool background_)
: OpenCLDeviceBase(info, stats, background_),
path_trace_program(this, "megakernel", "kernel.cl", "-D__COMPILE_ONLY_MEGAKERNEL__ ")
{
}
virtual void load_kernels(const DeviceRequestedFeatures& /*requested_features*/,
vector<OpenCLProgram*> &programs)
{
path_trace_program.add_kernel(ustring("path_trace"));
programs.push_back(&path_trace_program);
}
~OpenCLDeviceMegaKernel()
{
task_pool.stop();
path_trace_program.release();
}
void path_trace(RenderTile& rtile, int sample)
{
/* Cast arguments to cl types. */
cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
cl_mem d_buffer = CL_MEM_PTR(rtile.buffer);
cl_mem d_rng_state = CL_MEM_PTR(rtile.rng_state);
cl_int d_x = rtile.x;
cl_int d_y = rtile.y;
cl_int d_w = rtile.w;
cl_int d_h = rtile.h;
cl_int d_offset = rtile.offset;
cl_int d_stride = rtile.stride;
/* Sample arguments. */
cl_int d_sample = sample;
cl_kernel ckPathTraceKernel = path_trace_program(ustring("path_trace"));
cl_uint start_arg_index =
kernel_set_args(ckPathTraceKernel,
0,
d_data,
d_buffer,
d_rng_state);
#define KERNEL_TEX(type, ttype, name) \
set_kernel_arg_mem(ckPathTraceKernel, &start_arg_index, #name);
#include "kernel_textures.h"
#undef KERNEL_TEX
start_arg_index += kernel_set_args(ckPathTraceKernel,
start_arg_index,
d_sample,
d_x,
d_y,
d_w,
d_h,
d_offset,
d_stride);
enqueue_kernel(ckPathTraceKernel, d_w, d_h);
}
void thread_run(DeviceTask *task)
{
if(task->type == DeviceTask::FILM_CONVERT) {
film_convert(*task, task->buffer, task->rgba_byte, task->rgba_half);
}
else if(task->type == DeviceTask::SHADER) {
shader(*task);
}
else if(task->type == DeviceTask::PATH_TRACE) {
RenderTile tile;
/* Keep rendering tiles until done. */
while(task->acquire_tile(this, tile)) {
int start_sample = tile.start_sample;
int end_sample = tile.start_sample + tile.num_samples;
for(int sample = start_sample; sample < end_sample; sample++) {
if(task->get_cancel()) {
if(task->need_finish_queue == false)
break;
}
path_trace(tile, sample);
tile.sample = sample + 1;
task->update_progress(&tile);
}
/* Complete kernel execution before release tile */
/* This helps in multi-device render;
* The device that reaches the critical-section function
* release_tile waits (stalling other devices from entering
* release_tile) for all kernels to complete. If device1 (a
* slow-render device) reaches release_tile first then it would
* stall device2 (a fast-render device) from proceeding to render
* next tile.
*/
clFinish(cqCommandQueue);
task->release_tile(tile);
}
}
}
};
Device *opencl_create_mega_device(DeviceInfo& info, Stats& stats, bool background)
{
return new OpenCLDeviceMegaKernel(info, stats, background);
}
CCL_NAMESPACE_END
#endif

File diff suppressed because it is too large Load Diff

View File

@ -0,0 +1,761 @@
/*
* Copyright 2011-2013 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifdef WITH_OPENCL
#include "opencl.h"
#include "util_logging.h"
#include "util_path.h"
#include "util_time.h"
using std::cerr;
using std::endl;
CCL_NAMESPACE_BEGIN
OpenCLCache::Slot::ProgramEntry::ProgramEntry()
: program(NULL),
mutex(NULL)
{
}
OpenCLCache::Slot::ProgramEntry::ProgramEntry(const ProgramEntry& rhs)
: program(rhs.program),
mutex(NULL)
{
}
OpenCLCache::Slot::ProgramEntry::~ProgramEntry()
{
delete mutex;
}
OpenCLCache::Slot::Slot()
: context_mutex(NULL),
context(NULL)
{
}
OpenCLCache::Slot::Slot(const Slot& rhs)
: context_mutex(NULL),
context(NULL),
programs(rhs.programs)
{
}
OpenCLCache::Slot::~Slot()
{
delete context_mutex;
}
OpenCLCache& OpenCLCache::global_instance()
{
static OpenCLCache instance;
return instance;
}
cl_context OpenCLCache::get_context(cl_platform_id platform,
cl_device_id device,
thread_scoped_lock& slot_locker)
{
assert(platform != NULL);
OpenCLCache& self = global_instance();
thread_scoped_lock cache_lock(self.cache_lock);
pair<CacheMap::iterator,bool> ins = self.cache.insert(
CacheMap::value_type(PlatformDevicePair(platform, device), Slot()));
Slot &slot = ins.first->second;
/* create slot lock only while holding cache lock */
if(!slot.context_mutex)
slot.context_mutex = new thread_mutex;
/* need to unlock cache before locking slot, to allow store to complete */
cache_lock.unlock();
/* lock the slot */
slot_locker = thread_scoped_lock(*slot.context_mutex);
/* If the thing isn't cached */
if(slot.context == NULL) {
/* return with the caller's lock holder holding the slot lock */
return NULL;
}
/* the item was already cached, release the slot lock */
slot_locker.unlock();
cl_int ciErr = clRetainContext(slot.context);
assert(ciErr == CL_SUCCESS);
(void)ciErr;
return slot.context;
}
cl_program OpenCLCache::get_program(cl_platform_id platform,
cl_device_id device,
ustring key,
thread_scoped_lock& slot_locker)
{
assert(platform != NULL);
OpenCLCache& self = global_instance();
thread_scoped_lock cache_lock(self.cache_lock);
pair<CacheMap::iterator,bool> ins = self.cache.insert(
CacheMap::value_type(PlatformDevicePair(platform, device), Slot()));
Slot &slot = ins.first->second;
pair<Slot::EntryMap::iterator,bool> ins2 = slot.programs.insert(
Slot::EntryMap::value_type(key, Slot::ProgramEntry()));
Slot::ProgramEntry &entry = ins2.first->second;
/* create slot lock only while holding cache lock */
if(!entry.mutex)
entry.mutex = new thread_mutex;
/* need to unlock cache before locking slot, to allow store to complete */
cache_lock.unlock();
/* lock the slot */
slot_locker = thread_scoped_lock(*entry.mutex);
/* If the thing isn't cached */
if(entry.program == NULL) {
/* return with the caller's lock holder holding the slot lock */
return NULL;
}
/* the item was already cached, release the slot lock */
slot_locker.unlock();
cl_int ciErr = clRetainProgram(entry.program);
assert(ciErr == CL_SUCCESS);
(void)ciErr;
return entry.program;
}
void OpenCLCache::store_context(cl_platform_id platform,
cl_device_id device,
cl_context context,
thread_scoped_lock& slot_locker)
{
assert(platform != NULL);
assert(device != NULL);
assert(context != NULL);
OpenCLCache &self = global_instance();
thread_scoped_lock cache_lock(self.cache_lock);
CacheMap::iterator i = self.cache.find(PlatformDevicePair(platform, device));
cache_lock.unlock();
Slot &slot = i->second;
/* sanity check */
assert(i != self.cache.end());
assert(slot.context == NULL);
slot.context = context;
/* unlock the slot */
slot_locker.unlock();
/* increment reference count in OpenCL.
* The caller is going to release the object when done with it. */
cl_int ciErr = clRetainContext(context);
assert(ciErr == CL_SUCCESS);
(void)ciErr;
}
void OpenCLCache::store_program(cl_platform_id platform,
cl_device_id device,
cl_program program,
ustring key,
thread_scoped_lock& slot_locker)
{
assert(platform != NULL);
assert(device != NULL);
assert(program != NULL);
OpenCLCache &self = global_instance();
thread_scoped_lock cache_lock(self.cache_lock);
CacheMap::iterator i = self.cache.find(PlatformDevicePair(platform, device));
assert(i != self.cache.end());
Slot &slot = i->second;
Slot::EntryMap::iterator i2 = slot.programs.find(key);
assert(i2 != slot.programs.end());
Slot::ProgramEntry &entry = i2->second;
assert(entry.program == NULL);
cache_lock.unlock();
entry.program = program;
/* unlock the slot */
slot_locker.unlock();
/* Increment reference count in OpenCL.
* The caller is going to release the object when done with it.
*/
cl_int ciErr = clRetainProgram(program);
assert(ciErr == CL_SUCCESS);
(void)ciErr;
}
string OpenCLCache::get_kernel_md5()
{
OpenCLCache &self = global_instance();
thread_scoped_lock lock(self.kernel_md5_lock);
if(self.kernel_md5.empty()) {
self.kernel_md5 = path_files_md5_hash(path_get("kernel"));
}
return self.kernel_md5;
}
OpenCLDeviceBase::OpenCLProgram::OpenCLProgram(OpenCLDeviceBase *device, string program_name, string kernel_file, string kernel_build_options)
: device(device),
program_name(program_name),
kernel_file(kernel_file),
kernel_build_options(kernel_build_options)
{
loaded = false;
program = NULL;
}
OpenCLDeviceBase::OpenCLProgram::~OpenCLProgram()
{
release();
}
void OpenCLDeviceBase::OpenCLProgram::release()
{
for(map<ustring, cl_kernel>::iterator kernel = kernels.begin(); kernel != kernels.end(); ++kernel) {
if(kernel->second) {
clReleaseKernel(kernel->second);
kernel->second = NULL;
}
}
if(program) {
clReleaseProgram(program);
program = NULL;
}
}
void OpenCLDeviceBase::OpenCLProgram::add_kernel(ustring name)
{
if(!kernels.count(name)) {
kernels[name] = NULL;
}
}
bool OpenCLDeviceBase::OpenCLProgram::build_kernel(const string *debug_src)
{
string build_options;
build_options = device->kernel_build_options(debug_src) + kernel_build_options;
cl_int ciErr = clBuildProgram(program, 0, NULL, build_options.c_str(), NULL, NULL);
/* show warnings even if build is successful */
size_t ret_val_size = 0;
clGetProgramBuildInfo(program, device->cdDevice, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
if(ret_val_size > 1) {
vector<char> build_log(ret_val_size + 1);
clGetProgramBuildInfo(program, device->cdDevice, CL_PROGRAM_BUILD_LOG, ret_val_size, &build_log[0], NULL);
build_log[ret_val_size] = '\0';
/* Skip meaningless empty output from the NVidia compiler. */
if(!(ret_val_size == 2 && build_log[0] == '\n')) {
output_msg = string(&build_log[0]);
}
}
if(ciErr != CL_SUCCESS) {
error_msg = string("OpenCL build failed: ") + clewErrorString(ciErr);
return false;
}
return true;
}
bool OpenCLDeviceBase::OpenCLProgram::compile_kernel(const string *debug_src)
{
string source = "#include \"kernels/opencl/" + kernel_file + "\" // " + OpenCLCache::get_kernel_md5() + "\n";
/* We compile kernels consisting of many files. unfortunately OpenCL
* kernel caches do not seem to recognize changes in included files.
* so we force recompile on changes by adding the md5 hash of all files.
*/
source = path_source_replace_includes(source, path_get("kernel"));
if(debug_src) {
path_write_text(*debug_src, source);
}
size_t source_len = source.size();
const char *source_str = source.c_str();
cl_int ciErr;
program = clCreateProgramWithSource(device->cxContext,
1,
&source_str,
&source_len,
&ciErr);
if(ciErr != CL_SUCCESS) {
error_msg = string("OpenCL program creation failed: ") + clewErrorString(ciErr);
return false;
}
double starttime = time_dt();
log += "Build flags: " + kernel_build_options + "\n";
if(!build_kernel(debug_src))
return false;
log += "Kernel compilation of " + program_name + " finished in " + string_printf("%.2lfs.\n", time_dt() - starttime);
return true;
}
bool OpenCLDeviceBase::OpenCLProgram::load_binary(const string& clbin,
const string *debug_src)
{
/* read binary into memory */
vector<uint8_t> binary;
if(!path_read_binary(clbin, binary)) {
error_msg = "OpenCL failed to read cached binary " + clbin + ".";
return false;
}
/* create program */
cl_int status, ciErr;
size_t size = binary.size();
const uint8_t *bytes = &binary[0];
program = clCreateProgramWithBinary(device->cxContext, 1, &device->cdDevice,
&size, &bytes, &status, &ciErr);
if(status != CL_SUCCESS || ciErr != CL_SUCCESS) {
error_msg = "OpenCL failed create program from cached binary " + clbin + ": " + clewErrorString(status) + " " + clewErrorString(ciErr);
return false;
}
if(!build_kernel(debug_src))
return false;
return true;
}
bool OpenCLDeviceBase::OpenCLProgram::save_binary(const string& clbin)
{
size_t size = 0;
clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL);
if(!size)
return false;
vector<uint8_t> binary(size);
uint8_t *bytes = &binary[0];
clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(uint8_t*), &bytes, NULL);
return path_write_binary(clbin, binary);
}
void OpenCLDeviceBase::OpenCLProgram::load()
{
assert(device);
loaded = false;
string device_md5 = device->device_md5_hash(kernel_build_options);
/* Try to use cached kernel. */
thread_scoped_lock cache_locker;
ustring cache_key(program_name + device_md5);
program = device->load_cached_kernel(cache_key,
cache_locker);
if(!program) {
log += "OpenCL program " + program_name + " not found in cache.\n";
string basename = "cycles_kernel_" + program_name + "_" + device_md5 + "_" + OpenCLCache::get_kernel_md5();
basename = path_cache_get(path_join("kernels", basename));
string clbin = basename + ".clbin";
/* path to preprocessed source for debugging */
string clsrc, *debug_src = NULL;
if(OpenCLInfo::use_debug()) {
clsrc = basename + ".cl";
debug_src = &clsrc;
}
/* If binary kernel exists already, try use it. */
if(path_exists(clbin) && load_binary(clbin)) {
/* Kernel loaded from binary, nothing to do. */
log += "Loaded program from " + clbin + ".\n";
}
else {
log += "Kernel file " + clbin + " either doesn't exist or failed to be loaded by driver.\n";
/* If does not exist or loading binary failed, compile kernel. */
if(!compile_kernel(debug_src)) {
return;
}
/* Save binary for reuse. */
if(!save_binary(clbin)) {
log += "Saving compiled OpenCL kernel to " + clbin + " failed!";
}
}
/* Cache the program. */
device->store_cached_kernel(program,
cache_key,
cache_locker);
}
else {
log += "Found cached OpenCL program " + program_name + ".\n";
}
for(map<ustring, cl_kernel>::iterator kernel = kernels.begin(); kernel != kernels.end(); ++kernel) {
assert(kernel->second == NULL);
cl_int ciErr;
string name = "kernel_ocl_" + kernel->first.string();
kernel->second = clCreateKernel(program, name.c_str(), &ciErr);
if(device->opencl_error(ciErr)) {
error_msg = "Error getting kernel " + name + " from program " + program_name + ": " + clewErrorString(ciErr);
return;
}
}
loaded = true;
}
void OpenCLDeviceBase::OpenCLProgram::report_error()
{
if(loaded) return;
cerr << error_msg << endl;
if(!output_msg.empty()) {
cerr << "OpenCL kernel build output for " << program_name << ":" << endl;
cerr << output_msg << endl;
}
}
cl_kernel OpenCLDeviceBase::OpenCLProgram::operator()()
{
assert(kernels.size() == 1);
return kernels.begin()->second;
}
cl_kernel OpenCLDeviceBase::OpenCLProgram::operator()(ustring name)
{
assert(kernels.count(name));
return kernels[name];
}
cl_device_type OpenCLInfo::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;
}
}
bool OpenCLInfo::use_debug()
{
return DebugFlags().opencl.debug;
}
bool OpenCLInfo::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 OpenCLInfo::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 OpenCLInfo::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 OpenCLInfo::platform_version_check(cl_platform_id platform,
string *error)
{
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 OpenCLInfo::device_version_check(cl_device_id device,
string *error)
{
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 OpenCLInfo::get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices,
bool force_all)
{
const bool force_all_platforms = force_all ||
(DebugFlags().opencl.kernel_type != DebugFlags::OpenCL::KERNEL_DEFAULT);
const cl_device_type device_type = OpenCLInfo::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(!platform_version_check(platform_id)) {
FIRST_VLOG(2) << "Ignoring platform " << platform_name
<< " due to too old compiler version.";
continue;
}
num_devices = 0;
cl_int ciErr;
if((ciErr = 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: " << string(clewErrorString(ciErr));
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(!device_version_check(device_id)) {
FIRST_VLOG(2) << "Ignoring device " << device_name
<< " due to old compiler version.";
continue;
}
if(force_all_platforms ||
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;
}
CCL_NAMESPACE_END
#endif