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.

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();
 

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,

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;

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
 

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