Cycles code refactoring: change displace kernel into more generic shader

evaluate kernel, added background shader evaluate.

7 files changed, 62 insertions, 19 deletions

diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl
index 90eb7a2513f..479cf9b2e64 100644
--- a/intern/cycles/kernel/kernel.cl
+++ b/intern/cycles/kernel/kernel.cl
@@ -80,10 +80,10 @@ __kernel void kernel_ocl_tonemap(
 		kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
 }
 
-/*__kernel void kernel_ocl_displace(__global uint4 *input, __global float3 *offset, int sx)
+/*__kernel void kernel_ocl_shader(__global uint4 *input, __global float3 *output, int type, int sx)
 {
 	int x = sx + get_global_id(0);
 
-	kernel_displace(input, offset, x);
+	kernel_shader_evaluate(input, output, (ShaderEvalType)type, x);
 }*/
 
diff --git a/intern/cycles/kernel/kernel.cpp b/intern/cycles/kernel/kernel.cpp
index b4c3839dbd0..e66ddd86cd6 100644
--- a/intern/cycles/kernel/kernel.cpp
+++ b/intern/cycles/kernel/kernel.cpp
@@ -216,11 +216,11 @@ void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sam
 	kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
 }
 
-/* Displacement */
+/* Shader Evaluation */
 
-void kernel_cpu_displace(KernelGlobals *kg, uint4 *input, float3 *offset, int i)
+void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float3 *output, int type, int i)
 {
-	kernel_displace(kg, input, offset, i);
+	kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i);
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/kernel.cu b/intern/cycles/kernel/kernel.cu
index 71fc7ac3197..c97aeb67548 100644
--- a/intern/cycles/kernel/kernel.cu
+++ b/intern/cycles/kernel/kernel.cu
@@ -44,10 +44,10 @@ extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float4 *buffer, int
 		kernel_film_tonemap(NULL, rgba, buffer, sample, resolution, x, y, offset, stride);
 }
 
-extern "C" __global__ void kernel_cuda_displace(uint4 *input, float3 *offset, int sx)
+extern "C" __global__ void kernel_cuda_shader(uint4 *input, float3 *output, int type, int sx)
 {
 	int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
 
-	kernel_displace(NULL, input, offset, x);
+	kernel_shader_evaluate(NULL, input, output, (ShaderEvalType)type, x);
 }
 
diff --git a/intern/cycles/kernel/kernel.h b/intern/cycles/kernel/kernel.h
index 78247504b39..20d43c91169 100644
--- a/intern/cycles/kernel/kernel.h
+++ b/intern/cycles/kernel/kernel.h
@@ -40,14 +40,16 @@ void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_
 	int sample, int x, int y, int offset, int stride);
 void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer,
 	int sample, int resolution, int x, int y, int offset, int stride);
-void kernel_cpu_displace(KernelGlobals *kg, uint4 *input, float3 *offset, int i);
+void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float3 *output,
+	int type, int i);
 
 #ifdef WITH_OPTIMIZED_KERNEL
 void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state,
 	int sample, int x, int y, int offset, int stride);
 void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer,
 	int sample, int resolution, int x, int y, int offset, int stride);
-void kernel_cpu_optimized_displace(KernelGlobals *kg, uint4 *input, float3 *offset, int i);
+void kernel_cpu_optimized_shader(KernelGlobals *kg, uint4 *input, float3 *output,
+	int type, int i);
 #endif
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/kernel_displace.h b/intern/cycles/kernel/kernel_displace.h
index ef6c3810a75..c39e5e43dbb 100644
--- a/intern/cycles/kernel/kernel_displace.h
+++ b/intern/cycles/kernel/kernel_displace.h
@@ -18,17 +18,51 @@
 
 CCL_NAMESPACE_BEGIN
 
-__device void kernel_displace(KernelGlobals *kg, uint4 *input, float3 *offset, int i)
+__device void kernel_shader_evaluate(KernelGlobals *kg, uint4 *input, float3 *output, ShaderEvalType type, int i)
 {
-	/* setup shader data */
 	ShaderData sd;
 	uint4 in = input[i];
-	shader_setup_from_displace(kg, &sd, in.x, in.y, __int_as_float(in.z), __int_as_float(in.w));
+	float3 out;
 
-	/* evaluate */
-	float3 P = sd.P;
-	shader_eval_displacement(kg, &sd);
-	offset[i] = sd.P - P;
+	if(type == SHADER_EVAL_DISPLACE) {
+		/* setup shader data */
+		int object = in.x;
+		int prim = in.y;
+		float u = __int_as_float(in.z);
+		float v = __int_as_float(in.w);
+
+		shader_setup_from_displace(kg, &sd, object, prim, u, v);
+
+		/* evaluate */
+		float3 P = sd.P;
+		shader_eval_displacement(kg, &sd);
+		out = sd.P - P;
+	}
+	else { // SHADER_EVAL_BACKGROUND
+		/* setup ray */
+		Ray ray;
+
+		ray.P = make_float3(0.0f, 0.0f, 0.0f);
+		ray.D = make_float3(__int_as_float(in.x), __int_as_float(in.y), __int_as_float(in.z));
+		ray.t = 0.0f;
+
+#ifdef __RAY_DIFFERENTIALS__
+		ray.dD.dx = make_float3(0.0f, 0.0f, 0.0f);
+		ray.dD.dy = make_float3(0.0f, 0.0f, 0.0f);
+		ray.dP.dx = make_float3(0.0f, 0.0f, 0.0f);
+		ray.dP.dy = make_float3(0.0f, 0.0f, 0.0f);
+#endif
+
+		/* setup shader data */
+		shader_setup_from_background(kg, &sd, &ray);
+
+		/* evaluate */
+		int flag = 0; /* we can't know which type of BSDF this is for */
+		out = shader_eval_background(kg, &sd, flag);
+	}
+	
+	/* write output */
+	output[i] = out;
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/kernel_optimized.cpp b/intern/cycles/kernel/kernel_optimized.cpp
index ea43e01ab58..c437e06adfa 100644
--- a/intern/cycles/kernel/kernel_optimized.cpp
+++ b/intern/cycles/kernel/kernel_optimized.cpp
@@ -47,11 +47,11 @@ void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffe
 	kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
 }
 
-/* Displacement */
+/* Shader Evaluate */
 
-void kernel_cpu_optimized_displace(KernelGlobals *kg, uint4 *input, float3 *offset, int i)
+void kernel_cpu_optimized_shader(KernelGlobals *kg, uint4 *input, float3 *output, int type, int i)
 {
-	kernel_displace(kg, input, offset, i);
+	kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i);
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index 56db4d2b78a..2c03a34df1f 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -78,6 +78,13 @@ CCL_NAMESPACE_BEGIN
 //#define __MODIFY_TP__
 //#define __QBVH__
 
+/* Shader Evaluation */
+
+enum ShaderEvalType {
+	SHADER_EVAL_DISPLACE,
+	SHADER_EVAL_BACKGROUND
+};
+
 /* Path Tracing */
 
 enum PathTraceDimension {