Cycles OpenCL: make displacement and world importance sampling work.

1 files changed, 22 insertions, 3 deletions

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);
+}