[Bf-blender-cvs] [e7156b8] cycles_kernel_split: Change class hierarchy in device_opencl.cpp and refactor
varunsundar08
noreply at git.blender.org
Thu Apr 30 23:25:25 CEST 2015
Commit: e7156b828801a4b235ce65f477296b57385e49ba
Author: varunsundar08
Date: Thu Apr 30 11:10:59 2015 +0530
Branches: cycles_kernel_split
https://developer.blender.org/rBe7156b828801a4b235ce65f477296b57385e49ba
Change class hierarchy in device_opencl.cpp and refactor
Class hierarchy change :
The new class hierarchy is as follows,
OpenCLDeviceBase (inherits Device) - Contains error-handlers, texture functiions and other
common functions between MegaKernel and SplitKernel
OpenCLDeviceSplitKernel (inherits OpenCLDeviceBase) - Contains variables and functions specific to
split kernel implementation
OpenCLDeviceMegaKernel (inherits OpenCLDeviceBase) - Contains variables and functions specific to
megakernel
Refactor :
Change "if " to "if"
===================================================================
M intern/cycles/device/device_opencl.cpp
===================================================================
diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp
index 5e8b171..b1392fa 100644
--- a/intern/cycles/device/device_opencl.cpp
+++ b/intern/cycles/device/device_opencl.cpp
@@ -343,17 +343,15 @@ public:
}
};
-class OpenCLDevice : public Device
+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;
- cl_kernel ckPathTraceKernel;
cl_kernel ckShaderKernel;
cl_kernel ckBakeKernel;
cl_kernel ckFilmConvertByteKernel;
@@ -395,6 +393,7 @@ public:
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; \
@@ -415,8 +414,8 @@ public:
}
}
- OpenCLDevice(DeviceInfo& info, Stats &stats, bool background_)
- : Device(info, stats, background_)
+ OpenCLDeviceBase(DeviceInfo& info, Stats &stats, bool background_)
+ : Device(info, stats, background_)
{
cpPlatform = NULL;
cdDevice = NULL;
@@ -425,7 +424,6 @@ public:
null_mem = 0;
device_initialized = false;
- ckPathTraceKernel = NULL;
ckShaderKernel = NULL;
ckBakeKernel = NULL;
ckFilmConvertByteKernel = NULL;
@@ -455,7 +453,7 @@ public:
int num_base = 0;
int total_devices = 0;
- for(int platform = 0; platform < num_platforms; platform++) {
+ for (int platform = 0; platform < num_platforms; platform++) {
cl_uint num_devices;
if(opencl_error(clGetDeviceIDs(platforms[platform], opencl_device_type(), 0, NULL, &num_devices)))
@@ -532,7 +530,7 @@ public:
if(opencl_error(ciErr))
return;
- fprintf(stderr,"Device init succes\n");
+ fprintf(stderr, "Device init succes\n");
device_initialized = true;
}
@@ -578,114 +576,6 @@ public:
return true;
}
- bool load_binary(const string& kernel_path, const string& clbin, 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];
-
- cpProgram = 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(kernel_path, debug_src))
- return false;
-
- return true;
- }
-
- bool save_binary(const string& clbin)
- {
- size_t size = 0;
- clGetProgramInfo(cpProgram, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL);
-
- if(!size)
- return false;
-
- vector<uint8_t> binary(size);
- uint8_t *bytes = &binary[0];
-
- clGetProgramInfo(cpProgram, 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(const string& /*kernel_path*/, const string *debug_src = NULL)
- {
- string build_options = opencl_kernel_build_options(platform_name, debug_src);
-
- ciErr = clBuildProgram(cpProgram, 0, NULL, build_options.c_str(), NULL, NULL);
-
- /* show warnings even if build is successful */
- size_t ret_val_size = 0;
-
- clGetProgramBuildInfo(cpProgram, cdDevice, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
-
- if(ret_val_size > 1) {
- vector<char> build_log(ret_val_size+1);
- clGetProgramBuildInfo(cpProgram, cdDevice, CL_PROGRAM_BUILD_LOG, ret_val_size, &build_log[0], NULL);
-
- build_log[ret_val_size] = '\0';
- 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");
- return false;
- }
-
- return true;
- }
-
- bool compile_kernel(const string& kernel_path, const string& kernel_md5, 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 source = "#include \"kernel.cl\" // " + kernel_md5 + "\n";
- source = path_source_replace_includes(source, kernel_path);
-
- if(debug_src)
- path_write_text(*debug_src, source);
-
- size_t source_len = source.size();
- const char *source_str = source.c_str();
-
- cpProgram = clCreateProgramWithSource(cxContext, 1, &source_str, &source_len, &ciErr);
-
- if(opencl_error(ciErr))
- return false;
-
- double starttime = time_dt();
- printf("Compiling OpenCL kernel ...\n");
-
- if(!build_kernel(kernel_path, debug_src))
- return false;
-
- printf("Kernel compilation finished in %.2lfs.\n", time_dt() - starttime);
-
- return true;
- }
-
string device_md5_hash(string kernel_custom_build_option)
{
MD5Hash md5;
@@ -709,112 +599,31 @@ public:
return md5.get_hex();
}
- bool load_kernels(bool /*experimental*/)
- {
- /* 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 = OpenCLCache::get_program(cpPlatform, cdDevice, cache_locker);
-
- if(!cpProgram) {
- /* verify we have right opencl version */
- if(!opencl_version_check())
- return false;
-
- /* md5 hash to detect changes */
- string kernel_path = path_get("kernel");
- string kernel_md5 = path_files_md5_hash(kernel_path);
- string device_md5 = device_md5_hash("");
-
- /* path to cached binary */
- string clbin = string_printf("cycles_kernel_%s_%s.clbin", device_md5.c_str(), kernel_md5.c_str());
- clbin = path_user_get(path_join("cache", 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_user_get(path_join("cache", clsrc));
- debug_src = &clsrc;
- }
-
- /* if exists already, try use it */
- if(path_exists(clbin) && load_binary(kernel_path, clbin, debug_src)) {
- /* kernel loaded from binary */
- }
- else {
- /* if does not exist or loading binary failed, compile kernel */
- if(!compile_kernel(kernel_path, kernel_md5, debug_src))
- return false;
-
- /* save binary for reuse */
- if(!save_binary(clbin))
- return false;
- }
-
- /* cache the program */
- OpenCLCache::store_program(cpPlatform, cdDevice, cpProgram, cache_locker);
- }
-
- /* find kernels */
- ckPathTraceKernel = clCreateKernel(cpProgram, "kernel_ocl_path_trace", &ciErr);
- if(opencl_error(ciErr))
- return false;
-
- ckShaderKernel = clCreateKernel(cpProgram, "kernel_ocl_shader", &ciErr);
- if(opencl_error(ciErr))
- return false;
-
- ckBakeKernel = clCreateKernel(cpProgram, "kernel_ocl_bake", &ciErr);
- if (opencl_error(ciErr))
- return false;
-
- ckFilmConvertByteKernel = clCreateKernel(cpProgram, "kernel_ocl_convert_to_byte", &ciErr);
- if (opencl_error(ciErr))
- return false;
-
- ckFilmConvertHalfFloatKernel = clCreateKernel(cpProgram, "kernel_ocl_convert_to_half_float", &ciErr);
- if (opencl_error(ciErr))
- return false;
-
- return true;
- }
-
- ~OpenCLDevice()
+ ~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++) {
+ for (mt = const_mem_map.begin(); mt != const_mem_map.end(); mt++) {
mem_free(*(mt->second));
delete mt->second;
}
- if(ckPathTraceKernel)
- clReleaseKernel(ckPathTraceKernel);
-
- if (ckBakeKernel)
+ if(ckBakeKernel)
clReleaseKernel(ckBakeKernel);
- if (ckShaderKernel)
+ if(ckShaderKernel)
clReleaseKernel(ckShaderKernel);
- if (ckFilmConvertByteKernel)
+ if(ckFilmConvertByteKernel)
clReleaseKernel(ckFilmConvertByteKernel);
- if (ckFilmConvertHalfFloatKernel)
+ if(ckFilmConvertHalfFloatKernel)
clReleaseKernel(ckFilmConvertHalfFloatKernel);
- if (cpProgram)
+ if(cpProgram)
clReleaseProgram(cpProgram);
if(cqCommandQueue)
@@ -899,9 +708,9 @@ public:
}
void tex_alloc(const char *name,
- device_memory& mem,
- InterpolationType /*interpolation*/,
- bool /*periodic*/)
+ device_memory& mem,
+ InterpolationType /*interpolation*/,
+ bool /*periodic*/)
{
VLOG(1) << "Texture allocate: " << name << ", " << mem.memory_size() << " bytes.";
mem_alloc(mem, MEM_READ_ONLY);
@@ -927,7 +736,7 @@ public:
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);
+ return global_size + ((r == 0) ? 0 : group_size - r);
}
void enqueue_kernel(cl_kernel kernel, size_t w, size_t h)
@@ -937,61 +746,25 @@ public:
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);
+ 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};
+ 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[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
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list