Fix #29274: problem compiling cycles opencl kernel from directory with spaces.

Some drivers don't support passing include paths with spaces in them, nor does
the opencl spec specify anything about how to quote/escape such paths, so for
now we just resolved #includes ourselves. Alternative would have been to use c
preprocessor, but this also resolves all #ifdefs, which we do not want.

5 files changed, 47 insertions, 6 deletions

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<float> filter_table(FilterType type, float width)
 {
-	const int filter_table_size = FILTER_TABLE_SIZE;
+	const int filter_table_size = FILTER_TABLE_SIZE-1;
 	vector<float> filter_table_cdf(filter_table_size+1);
 	vector<float> 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<uint8_t>& binary)
 	return true;
 }
 
+static bool path_read_text(const string& path, string& text)
+{
+	vector<uint8_t> 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<uint8_t>& binary);
 bool path_read_binary(const string& path, vector<uint8_t>& binary);
 
+string path_source_replace_includes(const string& source, const string& path);
+
 CCL_NAMESPACE_END
 
 #endif