Merged revision(s) 57587-57670 from trunk/blender into soc-2013-dingto

12 files changed, 74 insertions, 30 deletions

diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index eda7fef67ef..184aaecc901 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -132,6 +132,12 @@ if(WITH_CYCLES_CUDA_BINARIES)
 	string(REGEX REPLACE ".*release ([0-9]+)\\.([0-9]+).*" "\\2" CUDA_VERSION_MINOR ${NVCC_OUT})
 	set(CUDA_VERSION "${CUDA_VERSION_MAJOR}${CUDA_VERSION_MINOR}")
 
+	# warn for other versions
+	if(CUDA_VERSION MATCHES "50")
+	else()
+		message(WARNING "CUDA version ${CUDA_VERSION_MAJOR}.${CUDA_VERSION_MINOR} detected, build may succeed but only CUDA 5.0 is officially supported")
+	endif()
+
 	# build for each arch
 	set(cuda_sources kernel.cu ${SRC_HEADERS} ${SRC_SVM_HEADERS} ${SRC_CLOSURE_HEADERS} ${SRC_UTIL_HEADERS})
 	set(cuda_cubins)
@@ -141,12 +147,6 @@ if(WITH_CYCLES_CUDA_BINARIES)
 
 		set(cuda_version_flags "-D__KERNEL_CUDA_VERSION__=${CUDA_VERSION}")
 
-		# warn for other versions
-		if(CUDA_VERSION MATCHES "50")
-		else()
-			message(STATUS "CUDA version ${CUDA_VERSION_MAJOR}.${CUDA_VERSION_MINOR} detected, build may succeed but only CUDA 5.0 is officially supported")
-		endif()
-
 		# build flags depending on CUDA version and arch
 		if(CUDA_VERSION LESS 50)
 			# CUDA 4.x
@@ -178,13 +178,17 @@ if(WITH_CYCLES_CUDA_BINARIES)
 			set(cuda_math_flags "--use_fast_math")
 		endif()
 		
-		add_custom_command(
-			OUTPUT ${cuda_cubin}
-			COMMAND ${CUDA_NVCC_EXECUTABLE} -arch=${arch} -m${CUDA_BITS} --cubin ${CMAKE_CURRENT_SOURCE_DIR}/kernel.cu -o ${CMAKE_CURRENT_BINARY_DIR}/${cuda_cubin} --ptxas-options="-v" ${cuda_arch_flags} ${cuda_version_flags} ${cuda_math_flags} -I${CMAKE_CURRENT_SOURCE_DIR}/../util -I${CMAKE_CURRENT_SOURCE_DIR}/svm -DCCL_NAMESPACE_BEGIN= -DCCL_NAMESPACE_END= -DNVCC
-			DEPENDS ${cuda_sources})
+		if(CUDA_VERSION LESS 50 AND ${arch} MATCHES "sm_35")
+			message(WARNING "Can't build kernel for CUDA sm_35 architecture, skipping")
+		else()
+			add_custom_command(
+				OUTPUT ${cuda_cubin}
+				COMMAND ${CUDA_NVCC_EXECUTABLE} -arch=${arch} -m${CUDA_BITS} --cubin ${CMAKE_CURRENT_SOURCE_DIR}/kernel.cu -o ${CMAKE_CURRENT_BINARY_DIR}/${cuda_cubin} --ptxas-options="-v" ${cuda_arch_flags} ${cuda_version_flags} ${cuda_math_flags} -I${CMAKE_CURRENT_SOURCE_DIR}/../util -I${CMAKE_CURRENT_SOURCE_DIR}/svm -DCCL_NAMESPACE_BEGIN= -DCCL_NAMESPACE_END= -DNVCC
+				DEPENDS ${cuda_sources})
 
-		delayed_install("${CMAKE_CURRENT_BINARY_DIR}" "${cuda_cubin}" ${CYCLES_INSTALL_PATH}/lib)
-		list(APPEND cuda_cubins ${cuda_cubin})
+			delayed_install("${CMAKE_CURRENT_BINARY_DIR}" "${cuda_cubin}" ${CYCLES_INSTALL_PATH}/lib)
+			list(APPEND cuda_cubins ${cuda_cubin})
+		endif()
 	endforeach()
 
 	add_custom_target(cycles_kernel_cuda ALL DEPENDS ${cuda_cubins})
diff --git a/intern/cycles/kernel/SConscript b/intern/cycles/kernel/SConscript
index 353ec1ce9d8..6459c3ed183 100644
--- a/intern/cycles/kernel/SConscript
+++ b/intern/cycles/kernel/SConscript
@@ -88,6 +88,10 @@ if env['WITH_BF_CYCLES_CUDA_BINARIES']:
 
 		# build flags depending on CUDA version and arch
         if cuda_version < 50:
+            if arch == "sm_35":
+                print("Can't build kernel for CUDA sm_35 architecture, skipping")
+                continue
+
             # CUDA 4.x
             if arch.startswith("sm_1"):
                 # sm_1x
diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl
index 22cb806b8e0..a745f5843fc 100644
--- a/intern/cycles/kernel/kernel.cl
+++ b/intern/cycles/kernel/kernel.cl
@@ -25,6 +25,7 @@
 
 #include "kernel_film.h"
 #include "kernel_path.h"
+#include "kernel_displace.h"
 
 __kernel void kernel_ocl_path_trace(
 	__constant KernelData *data,
@@ -80,10 +81,28 @@ __kernel void kernel_ocl_tonemap(
 		kernel_film_tonemap(kg, rgba, buffer, sample, x, y, offset, stride);
 }
 
-/*__kernel void kernel_ocl_shader(__global uint4 *input, __global float *output, int type, int sx)
+__kernel void kernel_ocl_shader(
+	__constant KernelData *data,
+	__global uint4 *input,
+	__global float4 *output,
+
+#define KERNEL_TEX(type, ttype, name) \
+	__global type *name,
+#include "kernel_textures.h"
+
+	int type, int sx, int sw)
 {
+	KernelGlobals kglobals, *kg = &kglobals;
+
+	kg->data = data;
+
+#define KERNEL_TEX(type, ttype, name) \
+	kg->name = name;
+#include "kernel_textures.h"
+
 	int x = sx + get_global_id(0);
 
-	kernel_shader_evaluate(input, output, (ShaderEvalType)type, x);
-}*/
+	if(x < sx + sw)
+		kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, x);
+}
 
diff --git a/intern/cycles/kernel/kernel_displace.h b/intern/cycles/kernel/kernel_displace.h
index c7fd03e7603..ae2e35e8d93 100644
--- a/intern/cycles/kernel/kernel_displace.h
+++ b/intern/cycles/kernel/kernel_displace.h
@@ -18,7 +18,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-__device void kernel_shader_evaluate(KernelGlobals *kg, uint4 *input, float4 *output, ShaderEvalType type, int i)
+__device void kernel_shader_evaluate(KernelGlobals *kg, __global uint4 *input, __global float4 *output, ShaderEvalType type, int i)
 {
 	ShaderData sd;
 	uint4 in = input[i];
diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h
index b895d1fcf52..6e1843df50d 100644
--- a/intern/cycles/kernel/kernel_path.h
+++ b/intern/cycles/kernel/kernel_path.h
@@ -249,7 +249,11 @@ __device float4 kernel_path_progressive(KernelGlobals *kg, RNG *rng, int sample,
 #endif
 	PathState state;
 	int rng_offset = PRNG_BASE_NUM;
+#ifdef __CMJ__
 	int num_samples = kernel_data.integrator.aa_samples;
+#else
+	int num_samples = 0;
+#endif
 
 	path_state_init(&state);
 
@@ -765,7 +769,11 @@ __device_noinline void kernel_path_non_progressive_lighting(KernelGlobals *kg, R
 	float min_ray_pdf, float ray_pdf, PathState state,
 	int rng_offset, PathRadiance *L, __global float *buffer)
 {
+#ifdef __CMJ__
 	int aa_samples = kernel_data.integrator.aa_samples;
+#else
+	int aa_samples = 0;
+#endif
 
 #ifdef __AO__
 	/* ambient occlusion */
@@ -964,7 +972,11 @@ __device float4 kernel_path_non_progressive(KernelGlobals *kg, RNG *rng, int sam
 	float ray_pdf = 0.0f;
 	PathState state;
 	int rng_offset = PRNG_BASE_NUM;
+#ifdef __CMJ__
 	int aa_samples = kernel_data.integrator.aa_samples;
+#else
+	int aa_samples = 0;
+#endif
 
 	path_state_init(&state);
 
@@ -1141,7 +1153,11 @@ __device void kernel_path_trace(KernelGlobals *kg,
 
 	float filter_u;
 	float filter_v;
+#ifdef __CMJ__
 	int num_samples = kernel_data.integrator.aa_samples;
+#else
+	int num_samples = 0;
+#endif
 
 	path_rng_init(kg, rng_state, sample, num_samples, &rng, x, y, &filter_u, &filter_v);
 
diff --git a/intern/cycles/kernel/kernel_random.h b/intern/cycles/kernel/kernel_random.h
index 20fc1fe2253..6292adff6a5 100644
--- a/intern/cycles/kernel/kernel_random.h
+++ b/intern/cycles/kernel/kernel_random.h
@@ -125,9 +125,9 @@ __device_inline float path_rng_1D(KernelGlobals *kg, RNG *rng, int sample, int n
 	float shift;
 
 	if(dimension & 1)
-		shift = (*rng >> 16)/((float)0xFFFF);
+		shift = (*rng >> 16) * (1.0f/(float)0xFFFF);
 	else
-		shift = (*rng & 0xFFFF)/((float)0xFFFF);
+		shift = (*rng & 0xFFFF) * (1.0f/(float)0xFFFF);
 
 	return r + shift - floorf(r + shift);
 #endif
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index c6b1e29ba26..348e8b8a5f1 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -157,10 +157,10 @@ CCL_NAMESPACE_BEGIN
 
 /* Shader Evaluation */
 
-enum ShaderEvalType {
+typedef enum ShaderEvalType {
 	SHADER_EVAL_DISPLACE,
 	SHADER_EVAL_BACKGROUND
-};
+} ShaderEvalType;
 
 /* Path Tracing
  * note we need to keep the u/v pairs at even values */
diff --git a/intern/cycles/kernel/shaders/node_bump.osl b/intern/cycles/kernel/shaders/node_bump.osl
index 6d0ec51a9da..321306b6c71 100644
--- a/intern/cycles/kernel/shaders/node_bump.osl
+++ b/intern/cycles/kernel/shaders/node_bump.osl
@@ -52,6 +52,6 @@ surface node_bump(
 	
 	/* compute and output perturbed normal */
 	NormalOut = normalize(absdet * NormalIn - dist * sign(det) * surfgrad);
-	NormalOut = normalize(strength*NormalOut + (1.0 - strength)*NormalIn);
+	NormalOut = normalize(strength * NormalOut + (1.0 - strength) * NormalIn);
 }
 
diff --git a/intern/cycles/kernel/shaders/node_color.h b/intern/cycles/kernel/shaders/node_color.h
index 92bd39b5828..7fdef65f7f0 100644
--- a/intern/cycles/kernel/shaders/node_color.h
+++ b/intern/cycles/kernel/shaders/node_color.h
@@ -50,8 +50,8 @@ color color_scene_linear_to_srgb(color c)
 
 color color_unpremultiply(color c, float alpha)
 {
-	if(alpha != 1.0 && alpha != 0.0)
-		return c/alpha;
+	if (alpha != 1.0 && alpha != 0.0)
+		return c / alpha;
 
 	return c;
 }
diff --git a/intern/cycles/kernel/shaders/node_environment_texture.osl b/intern/cycles/kernel/shaders/node_environment_texture.osl
index bddef418c3d..230116d3d77 100644
--- a/intern/cycles/kernel/shaders/node_environment_texture.osl
+++ b/intern/cycles/kernel/shaders/node_environment_texture.osl
@@ -49,6 +49,7 @@ shader node_environment_texture(
 	string projection = "Equirectangular",
 	string color_space = "sRGB",
 	int is_float = 1,
+	int use_alpha = 1,
 	output color Color = 0.0,
 	output float Alpha = 1.0)
 {
@@ -67,7 +68,7 @@ shader node_environment_texture(
 	/* todo: use environment for better texture filtering of equirectangular */
 	Color = (color)texture(filename, p[0], 1.0 - p[1], "wrap", "periodic", "alpha", Alpha);
 
-	if (isconnected(Alpha)) {
+	if (use_alpha) {
 		Color = color_unpremultiply(Color, Alpha);
 
 		if (!is_float)
diff --git a/intern/cycles/kernel/shaders/node_image_texture.osl b/intern/cycles/kernel/shaders/node_image_texture.osl
index 6ccc7ebc651..92c625e99c4 100644
--- a/intern/cycles/kernel/shaders/node_image_texture.osl
+++ b/intern/cycles/kernel/shaders/node_image_texture.osl
@@ -30,8 +30,9 @@ color image_texture_lookup(string filename, string color_space, float u, float v
 			rgb = min(rgb, 1.0);
 	}
 
-	if (color_space == "sRGB")
+	if (color_space == "sRGB") {
 		rgb = color_srgb_to_scene_linear(rgb);
+	}
 
 	return rgb;
 }
@@ -45,6 +46,7 @@ shader node_image_texture(
 	string projection = "Flat",
 	float projection_blend = 0.0,
 	int is_float = 1,
+	int use_alpha = 1,
 	output color Color = 0.0,
 	output float Alpha = 1.0)
 {
@@ -53,8 +55,6 @@ shader node_image_texture(
 	if (use_mapping)
 		p = transform(mapping, p);
 	
-	int use_alpha = isconnected(Alpha);
-
 	if (projection == "Flat") {
 		Color = image_texture_lookup(filename, color_space, p[0], p[1], Alpha, use_alpha, is_float);
 	}
diff --git a/intern/cycles/kernel/shaders/node_texture.h b/intern/cycles/kernel/shaders/node_texture.h
index 0a73c0907d9..645be23c6a5 100644
--- a/intern/cycles/kernel/shaders/node_texture.h
+++ b/intern/cycles/kernel/shaders/node_texture.h
@@ -162,7 +162,7 @@ float safe_noise(point p, int type)
 		f = noise(p);
 
 	/* can happen for big coordinates, things even out to 0.5 then anyway */
-	if(!isfinite(f))
+	if (!isfinite(f))
 		return 0.5;
 	
 	return f;
@@ -254,12 +254,12 @@ float noise_turbulence(point p, string basis, float details, int hard)
 		if (hard)
 			t = fabs(2.0 * t - 1.0);
 
-		float sum2 = sum + t*amp;
+		float sum2 = sum + t * amp;
 
 		sum *= ((float)(1 << n) / (float)((1 << (n + 1)) - 1));
 		sum2 *= ((float)(1 << (n + 1)) / (float)((1 << (n + 2)) - 1));
 
-		return (1.0 - rmd)*sum + rmd*sum2;
+		return (1.0 - rmd) * sum + rmd * sum2;
 	}
 	else {
 		sum *= ((float)(1 << n) / (float)((1 << (n + 1)) - 1));