[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [42071] trunk/blender/intern/cycles: Fix #29274: problem compiling cycles opencl kernel from directory with spaces.

[Brecht Van Lommel]

Revision: 42071
          http://projects.blender.org/scm/viewvc.php?view=rev&root=bf-blender&revision=42071
Author:   blendix
Date:     2011-11-22 16:38:58 +0000 (Tue, 22 Nov 2011)
Log Message:
-----------
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.

Modified Paths:
--------------
    trunk/blender/intern/cycles/device/device_opencl.cpp
    trunk/blender/intern/cycles/kernel/kernel.cl
    trunk/blender/intern/cycles/render/filter.cpp
    trunk/blender/intern/cycles/util/util_path.cpp
    trunk/blender/intern/cycles/util/util_path.h

Modified: trunk/blender/intern/cycles/device/device_opencl.cpp
===================================================================
--- trunk/blender/intern/cycles/device/device_opencl.cpp	2011-11-22 15:33:44 UTC (rev 42070)
+++ trunk/blender/intern/cycles/device/device_opencl.cpp	2011-11-22 16:38:58 UTC (rev 42071)
@@ -278,10 +278,7 @@
 
 	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 @@
 		   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();
 

Modified: trunk/blender/intern/cycles/kernel/kernel.cl
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel.cl	2011-11-22 15:33:44 UTC (rev 42070)
+++ trunk/blender/intern/cycles/kernel/kernel.cl	2011-11-22 16:38:58 UTC (rev 42071)
@@ -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,

Modified: trunk/blender/intern/cycles/render/filter.cpp
===================================================================
--- trunk/blender/intern/cycles/render/filter.cpp	2011-11-22 15:33:44 UTC (rev 42070)
+++ trunk/blender/intern/cycles/render/filter.cpp	2011-11-22 16:38:58 UTC (rev 42071)
@@ -53,7 +53,7 @@
 
 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;

Modified: trunk/blender/intern/cycles/util/util_path.cpp
===================================================================
--- trunk/blender/intern/cycles/util/util_path.cpp	2011-11-22 15:33:44 UTC (rev 42070)
+++ trunk/blender/intern/cycles/util/util_path.cpp	2011-11-22 16:38:58 UTC (rev 42071)
@@ -162,5 +162,46 @@
 	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
 

Modified: trunk/blender/intern/cycles/util/util_path.h
===================================================================
--- trunk/blender/intern/cycles/util/util_path.h	2011-11-22 15:33:44 UTC (rev 42070)
+++ trunk/blender/intern/cycles/util/util_path.h	2011-11-22 16:38:58 UTC (rev 42071)
@@ -46,6 +46,8 @@
 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