diff options
Diffstat (limited to 'intern/cycles/kernel')
-rw-r--r-- | intern/cycles/kernel/kernel.cl | 25 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_displace.h | 2 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_types.h | 4 |
3 files changed, 25 insertions, 6 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); +} 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_types.h b/intern/cycles/kernel/kernel_types.h index a9f1831fc7d..3bd0d5c3561 100644 --- a/intern/cycles/kernel/kernel_types.h +++ b/intern/cycles/kernel/kernel_types.h @@ -150,10 +150,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 */ |