Cr8-xyz/blender
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

Cycles: Prepare code for OpenCL camera/motion blur

Browse Source 
The kernels are now compiling just fine, but there're some issues
during rendering. This is still to be investigated.
This commit is contained in:
Sergey Sharybin
2015-05-14 18:46:26 +05:00
parent 5a63edb929
commit 3d3d805b64
5 changed files with 19 additions and 18 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
intern/cycles/kernel/geom/geom_motion_triangle.h
View File

@@ -134,7 +134,7 @@ ccl_device_inline float3 motion_triangle_refine(KernelGlobals *kg, ShaderData *s
return P; return P;
} }
#ifdef __OBJECT_MOTION__ #ifdef __OBJECT_MOTION__
Transform tfm = sd->ob_itfm; Transform tfm = ccl_fetch(sd, ob_itfm);
#else #else
Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_INVERSE_TRANSFORM); Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_INVERSE_TRANSFORM);
#endif #endif
@@ -161,7 +161,7 @@ ccl_device_inline float3 motion_triangle_refine(KernelGlobals *kg, ShaderData *s
if(isect->object != OBJECT_NONE) { if(isect->object != OBJECT_NONE) {
#ifdef __OBJECT_MOTION__ #ifdef __OBJECT_MOTION__
Transform tfm = sd->ob_tfm; Transform tfm = ccl_fetch(sd, ob_tfm);
#else #else
Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_TRANSFORM); Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_TRANSFORM);
#endif #endif
@@ -187,7 +187,7 @@ ccl_device_inline float3 motion_triangle_refine_subsurface(KernelGlobals *kg, Sh
#ifdef __INTERSECTION_REFINE__ #ifdef __INTERSECTION_REFINE__
if(isect->object != OBJECT_NONE) { if(isect->object != OBJECT_NONE) {
#ifdef __OBJECT_MOTION__ #ifdef __OBJECT_MOTION__
Transform tfm = sd->ob_itfm; Transform tfm = ccl_fetch(sd, ob_itfm);
#else #else
Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_INVERSE_TRANSFORM); Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_INVERSE_TRANSFORM);
#endif #endif
@@ -213,7 +213,7 @@ ccl_device_inline float3 motion_triangle_refine_subsurface(KernelGlobals *kg, Sh
if(isect->object != OBJECT_NONE) { if(isect->object != OBJECT_NONE) {
#ifdef __OBJECT_MOTION__ #ifdef __OBJECT_MOTION__
Transform tfm = sd->ob_tfm; Transform tfm = ccl_fetch(sd, ob_tfm);
#else #else
Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_TRANSFORM); Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_TRANSFORM);
#endif #endif

4
intern/cycles/kernel/geom/geom_object.h
View File

@@ -453,7 +453,7 @@ ccl_device_inline void bvh_instance_pop_factor(KernelGlobals *kg, int object, co
#ifdef __OBJECT_MOTION__ #ifdef __OBJECT_MOTION__
/* Transform ray into object space to enter motion blurred object in BVH */ /* Transform ray into object space to enter motion blurred object in BVH */
ccl_device_inline void bvh_instance_motion_push(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *dir, float3 *idir, float *t, Transform *tfm) ccl_device_inline void bvh_instance_motion_push(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *dir, float3 *idir, ccl_addr_space float *t, Transform *tfm)
{ {
Transform itfm; Transform itfm;
*tfm = object_fetch_transform_motion_test(kg, object, ray->time, &itfm); *tfm = object_fetch_transform_motion_test(kg, object, ray->time, &itfm);
@@ -497,7 +497,7 @@ ccl_device_inline void qbvh_instance_motion_push(KernelGlobals *kg, int object,
/* Transorm ray to exit motion blurred object in BVH */ /* Transorm ray to exit motion blurred object in BVH */
ccl_device_inline void bvh_instance_motion_pop(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *dir, float3 *idir, float *t, Transform *tfm) ccl_device_inline void bvh_instance_motion_pop(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *dir, float3 *idir, ccl_addr_space float *t, Transform *tfm)
{ {
if(*t != FLT_MAX) if(*t != FLT_MAX)
*t *= len(transform_direction(tfm, 1.0f/(*idir))); *t *= len(transform_direction(tfm, 1.0f/(*idir)));

8
intern/cycles/kernel/geom/geom_triangle_intersect.h
View File

@@ -306,7 +306,7 @@ ccl_device_inline float3 triangle_refine(KernelGlobals *kg,
return P; return P;
} }
#ifdef __OBJECT_MOTION__ #ifdef __OBJECT_MOTION__
Transform tfm = sd->ob_itfm; Transform tfm = ccl_fetch(sd, ob_itfm);
#else #else
Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_INVERSE_TRANSFORM); Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_INVERSE_TRANSFORM);
#endif #endif
@@ -332,7 +332,7 @@ ccl_device_inline float3 triangle_refine(KernelGlobals *kg,
if(isect->object != OBJECT_NONE) { if(isect->object != OBJECT_NONE) {
#ifdef __OBJECT_MOTION__ #ifdef __OBJECT_MOTION__
Transform tfm = sd->ob_tfm; Transform tfm = ccl_fetch(sd, ob_tfm);
#else #else
Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_TRANSFORM); Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_TRANSFORM);
#endif #endif
@@ -361,7 +361,7 @@ ccl_device_inline float3 triangle_refine_subsurface(KernelGlobals *kg,
#ifdef __INTERSECTION_REFINE__ #ifdef __INTERSECTION_REFINE__
if(isect->object != OBJECT_NONE) { if(isect->object != OBJECT_NONE) {
#ifdef __OBJECT_MOTION__ #ifdef __OBJECT_MOTION__
Transform tfm = sd->ob_itfm; Transform tfm = ccl_fetch(sd, ob_itfm);
#else #else
Transform tfm = object_fetch_transform(kg, Transform tfm = object_fetch_transform(kg,
isect->object, isect->object,
@@ -389,7 +389,7 @@ ccl_device_inline float3 triangle_refine_subsurface(KernelGlobals *kg,
if(isect->object != OBJECT_NONE) { if(isect->object != OBJECT_NONE) {
#ifdef __OBJECT_MOTION__ #ifdef __OBJECT_MOTION__
Transform tfm = sd->ob_tfm; Transform tfm = ccl_fetch(sd, ob_tfm);
#else #else
Transform tfm = object_fetch_transform(kg, Transform tfm = object_fetch_transform(kg,
isect->object, isect->object,

7
intern/cycles/kernel/kernel_direct_lighting.cl
View File

@@ -105,13 +105,14 @@ __kernel void kernel_ocl_path_trace_direct_lighting(
float light_u, light_v; float light_u, light_v;
path_state_rng_2D(kg, rng, state, PRNG_LIGHT_U, &light_u, &light_v); path_state_rng_2D(kg, rng, state, PRNG_LIGHT_U, &light_u, &light_v);
#ifdef __OBJECT_MOTION__
light_ray.time = sd->time;
#endif
LightSample ls; LightSample ls;
light_sample(kg, light_t, light_u, light_v, ccl_fetch(sd, time), ccl_fetch(sd, P), state->bounce, &ls); light_sample(kg, light_t, light_u, light_v, ccl_fetch(sd, time), ccl_fetch(sd, P), state->bounce, &ls);
Ray light_ray; Ray light_ray;
#ifdef __OBJECT_MOTION__
light_ray.time = ccl_fetch(sd, time);
#endif
BsdfEval L_light; BsdfEval L_light;
bool is_lamp; bool is_lamp;
if(direct_emission(kg, sd, &ls, &light_ray, &L_light, &is_lamp, state->bounce, state->transparent_bounce, sd_DL)) { if(direct_emission(kg, sd, &ls, &light_ray, &L_light, &is_lamp, state->bounce, state->transparent_bounce, sd_DL)) {

10
intern/cycles/kernel/kernel_shader.h
View File

@@ -37,13 +37,13 @@ CCL_NAMESPACE_BEGIN
#ifdef __OBJECT_MOTION__ #ifdef __OBJECT_MOTION__
ccl_device void shader_setup_object_transforms(KernelGlobals *kg, ShaderData *sd, float time) ccl_device void shader_setup_object_transforms(KernelGlobals *kg, ShaderData *sd, float time)
{ {
if(sd->flag & SD_OBJECT_MOTION) { if(ccl_fetch(sd, flag) & SD_OBJECT_MOTION) {
sd->ob_tfm = object_fetch_transform_motion(kg, sd->object, time); ccl_fetch(sd, ob_tfm) = object_fetch_transform_motion(kg, ccl_fetch(sd, object), time);
sd->ob_itfm = transform_quick_inverse(sd->ob_tfm); ccl_fetch(sd, ob_itfm) = transform_quick_inverse(ccl_fetch(sd, ob_tfm));
} }
else { else {
sd->ob_tfm = object_fetch_transform(kg, sd->object, OBJECT_TRANSFORM); ccl_fetch(sd, ob_tfm) = object_fetch_transform(kg, ccl_fetch(sd, object), OBJECT_TRANSFORM);
sd->ob_itfm = object_fetch_transform(kg, sd->object, OBJECT_INVERSE_TRANSFORM); ccl_fetch(sd, ob_itfm) = object_fetch_transform(kg, ccl_fetch(sd, object), OBJECT_INVERSE_TRANSFORM);
} }
} }
#endif #endif