diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp index f75928c1b80..e0b4f67da05 100644 --- a/intern/cycles/device/device_opencl.cpp +++ b/intern/cycles/device/device_opencl.cpp @@ -278,10 +278,7 @@ public: bool build_kernel(const string& kernel_path) { - string build_options = ""; - - build_options += "-I " + kernel_path + ""; /* todo: escape path, but it doesn't get parsed correct? */ - build_options += kernel_build_options(); + string build_options = kernel_build_options(); ciErr = clBuildProgram(cpProgram, 0, NULL, build_options.c_str(), NULL, NULL); @@ -312,6 +309,8 @@ public: 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); + size_t source_len = source.size(); const char *source_str = source.c_str(); diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl index 68ca24af58e..c00bc3fe957 100644 --- a/intern/cycles/kernel/kernel.cl +++ b/intern/cycles/kernel/kernel.cl @@ -25,7 +25,6 @@ #include "kernel_film.h" #include "kernel_path.h" -//#include "kernel_displace.h" __kernel void kernel_ocl_path_trace( __constant KernelData *data, diff --git a/intern/cycles/render/filter.cpp b/intern/cycles/render/filter.cpp index c000f1a0636..9bcf57b5a27 100644 --- a/intern/cycles/render/filter.cpp +++ b/intern/cycles/render/filter.cpp @@ -53,7 +53,7 @@ static float filter_func_gaussian(float v, float width) static vector filter_table(FilterType type, float width) { - const int filter_table_size = FILTER_TABLE_SIZE; + const int filter_table_size = FILTER_TABLE_SIZE-1; vector filter_table_cdf(filter_table_size+1); vector filter_table(filter_table_size+1); float (*filter_func)(float, float) = NULL; diff --git a/intern/cycles/util/util_path.cpp b/intern/cycles/util/util_path.cpp index 55e6a95b54a..69069a3bbce 100644 --- a/intern/cycles/util/util_path.cpp +++ b/intern/cycles/util/util_path.cpp @@ -162,5 +162,46 @@ bool path_read_binary(const string& path, vector& binary) return true; } +static bool path_read_text(const string& path, string& text) +{ + vector binary; + + if(!path_exists(path) || !path_read_binary(path, binary)) + return false; + + const char *str = (const char*)&binary[0]; + size_t size = binary.size(); + text = string(str, size); + + return true; +} + +string path_source_replace_includes(const string& source_, const string& path) +{ + /* our own little c preprocessor that replaces #includes with the file + contents, to work around issue of opencl drivers not supporting + include paths with spaces in them */ + string source = source_; + const string include = "#include \""; + size_t n, pos = 0; + + while((n = source.find(include, pos)) != string::npos) { + size_t n_start = n + include.size(); + size_t n_end = source.find("\"", n_start); + string filename = source.substr(n_start, n_end - n_start); + + string text, filepath = path_join(path, filename); + + if(path_read_text(filepath, text)) { + text = path_source_replace_includes(text, path_dirname(filepath)); + source.replace(n, n_end + 1 - n, "\n" + text + "\n"); + } + else + pos = n_end; + } + + return source; +} + CCL_NAMESPACE_END diff --git a/intern/cycles/util/util_path.h b/intern/cycles/util/util_path.h index 4a9d45ec594..6cba6a3158f 100644 --- a/intern/cycles/util/util_path.h +++ b/intern/cycles/util/util_path.h @@ -46,6 +46,8 @@ void path_create_directories(const string& path); bool path_write_binary(const string& path, const vector& binary); bool path_read_binary(const string& path, vector& binary); +string path_source_replace_includes(const string& source, const string& path); + CCL_NAMESPACE_END #endif