[Bf-blender-cvs] [c70fdb3] cycles_kernel_split: Cycles kernel split: Move address space specification to Intersection typedef
Sergey Sharybin
noreply at git.blender.org
Thu May 7 11:57:45 CEST 2015
Commit: c70fdb3f7380ffb4e0522ab021f60ac64d345935
Author: Sergey Sharybin
Date: Thu May 7 14:16:46 2015 +0500
Branches: cycles_kernel_split
https://developer.blender.org/rBc70fdb3f7380ffb4e0522ab021f60ac64d345935
Cycles kernel split: Move address space specification to Intersection typedef
===================================================================
M intern/cycles/kernel/geom/geom_bvh.h
M intern/cycles/kernel/geom/geom_bvh_traversal.h
M intern/cycles/kernel/geom/geom_motion_triangle.h
M intern/cycles/kernel/geom/geom_triangle_intersect.h
M intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl
M intern/cycles/kernel/kernel_LampEmission.cl
M intern/cycles/kernel/kernel_SceneIntersect.cl
M intern/cycles/kernel/kernel_ShaderEval.cl
M intern/cycles/kernel/kernel_ShadowBlocked.cl
M intern/cycles/kernel/kernel_shader.h
M intern/cycles/kernel/kernel_shadow.h
M intern/cycles/kernel/kernel_types.h
===================================================================
diff --git a/intern/cycles/kernel/geom/geom_bvh.h b/intern/cycles/kernel/geom/geom_bvh.h
index a15b7a8..3d0d406 100644
--- a/intern/cycles/kernel/geom/geom_bvh.h
+++ b/intern/cycles/kernel/geom/geom_bvh.h
@@ -216,7 +216,7 @@ CCL_NAMESPACE_BEGIN
#undef BVH_NAME_EVAL
#undef BVH_FUNCTION_FULL_NAME
-ccl_device_intersect bool scene_intersect(KernelGlobals *kg, const Ray *ray, const uint visibility, ccl_addr_space Intersection *isect,
+ccl_device_intersect bool scene_intersect(KernelGlobals *kg, const Ray *ray, const uint visibility, Intersection *isect,
uint *lcg_state, float difl, float extmax)
{
#ifdef __OBJECT_MOTION__
@@ -447,6 +447,7 @@ ccl_device_inline float3 ray_offset(float3 P, float3 Ng)
#endif
}
+#if defined(__SHADOW_RECORD_ALL__) || defined (__VOLUME_RECORD_ALL__)
/* ToDo: Move to another file? */
ccl_device int intersections_compare(const void *a, const void *b)
{
@@ -460,6 +461,7 @@ ccl_device int intersections_compare(const void *a, const void *b)
else
return 0;
}
+#endif
CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/geom/geom_bvh_traversal.h b/intern/cycles/kernel/geom/geom_bvh_traversal.h
index bc2483e..87f8ccd 100644
--- a/intern/cycles/kernel/geom/geom_bvh_traversal.h
+++ b/intern/cycles/kernel/geom/geom_bvh_traversal.h
@@ -34,7 +34,7 @@
ccl_device bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
const Ray *ray,
- ccl_addr_space Intersection *isect,
+ Intersection *isect,
const uint visibility
#if BVH_FEATURE(BVH_HAIR_MINIMUM_WIDTH)
, uint *lcg_state,
@@ -401,7 +401,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
ccl_device_inline bool BVH_FUNCTION_NAME(KernelGlobals *kg,
const Ray *ray,
- ccl_addr_space Intersection *isect,
+ Intersection *isect,
const uint visibility
#if BVH_FEATURE(BVH_HAIR_MINIMUM_WIDTH)
, uint *lcg_state,
diff --git a/intern/cycles/kernel/geom/geom_motion_triangle.h b/intern/cycles/kernel/geom/geom_motion_triangle.h
index f26ca96..47057ea 100644
--- a/intern/cycles/kernel/geom/geom_motion_triangle.h
+++ b/intern/cycles/kernel/geom/geom_motion_triangle.h
@@ -122,7 +122,7 @@ ccl_device_inline void motion_triangle_vertices(KernelGlobals *kg, int object, i
* far the precision is often not so good, this reintersects the primitive from
* a closer distance. */
-ccl_device_inline float3 motion_triangle_refine(KernelGlobals *kg, ShaderData *sd, const ccl_addr_space Intersection *isect, const Ray *ray, float3 verts[3])
+ccl_device_inline float3 motion_triangle_refine(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray, float3 verts[3])
{
float3 P = ray->P;
float3 D = ray->D;
@@ -233,7 +233,7 @@ ccl_device_inline float3 motion_triangle_refine_subsurface(KernelGlobals *kg, Sh
* normals */
/* return 3 triangle vertex normals */
-ccl_device_noinline void motion_triangle_shader_setup(KernelGlobals *kg, ShaderData *sd, const ccl_addr_space Intersection *isect, const Ray *ray, bool subsurface)
+ccl_device_noinline void motion_triangle_shader_setup(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray, bool subsurface)
{
/* get shader */
sd_fetch(shader) = kernel_tex_fetch(__tri_shader, sd_fetch(prim));
diff --git a/intern/cycles/kernel/geom/geom_triangle_intersect.h b/intern/cycles/kernel/geom/geom_triangle_intersect.h
index 8fc7014..3990bae 100644
--- a/intern/cycles/kernel/geom/geom_triangle_intersect.h
+++ b/intern/cycles/kernel/geom/geom_triangle_intersect.h
@@ -96,7 +96,7 @@ ccl_device_inline float xor_signmast(float x, int y)
ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
const IsectPrecalc *isect_precalc,
- ccl_addr_space Intersection *isect,
+ Intersection *isect,
float3 P,
uint visibility,
int object,
@@ -293,7 +293,7 @@ ccl_device_inline void triangle_intersect_subsurface(
ccl_device_inline float3 triangle_refine(KernelGlobals *kg,
ShaderData *sd,
- const ccl_addr_space Intersection *isect,
+ const Intersection *isect,
const Ray *ray)
{
float3 P = ray->P;
diff --git a/intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl b/intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl
index 2489d61..693211d 100644
--- a/intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl
+++ b/intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl
@@ -82,7 +82,7 @@ __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_AO
ccl_global float *L_transparent_coop, /* Required for handling holdout material */
ccl_global PathRadiance *PathRadiance_coop, /* Required for "kernel_write_data_passes" and indirect primitive emission */
ccl_global PathState *PathState_coop, /* Required throughout the kernel and AO */
- ccl_global Intersection *Intersection_coop, /* Required for indirect primitive emission */
+ Intersection *Intersection_coop, /* Required for indirect primitive emission */
ccl_global float3 *AOAlpha_coop, /* Required for AO */
ccl_global float3 *AOBSDF_coop, /* Required for AO */
ccl_global Ray *AOLightRay_coop, /* Required for AO */
diff --git a/intern/cycles/kernel/kernel_LampEmission.cl b/intern/cycles/kernel/kernel_LampEmission.cl
index 23a0fbe..26e2ade 100644
--- a/intern/cycles/kernel/kernel_LampEmission.cl
+++ b/intern/cycles/kernel/kernel_LampEmission.cl
@@ -48,7 +48,7 @@ __kernel void kernel_ocl_path_trace_LampEmission_SPLIT_KERNEL(
ccl_global PathRadiance *PathRadiance_coop, /* Required for lamp emission */
ccl_global Ray *Ray_coop, /* Required for lamp emission */
ccl_global PathState *PathState_coop, /* Required for lamp emission */
- ccl_global Intersection *Intersection_coop, /* Required for lamp emission */
+ Intersection *Intersection_coop, /* Required for lamp emission */
ccl_global char *ray_state, /* Denotes the state of each ray */
int sw, int sh,
ccl_global int *Queue_data, /* Memory for queues */
diff --git a/intern/cycles/kernel/kernel_SceneIntersect.cl b/intern/cycles/kernel/kernel_SceneIntersect.cl
index ff64782..34bb359 100644
--- a/intern/cycles/kernel/kernel_SceneIntersect.cl
+++ b/intern/cycles/kernel/kernel_SceneIntersect.cl
@@ -69,7 +69,7 @@ __kernel void kernel_ocl_path_trace_SceneIntersect_SPLIT_KERNEL(
ccl_global uint *rng_coop,
ccl_global Ray *Ray_coop, /* Required for scene_intersect */
ccl_global PathState *PathState_coop, /* Required for scene_intersect */
- ccl_global Intersection *Intersection_coop, /* Required for scene_intersect */
+ Intersection *Intersection_coop, /* Required for scene_intersect */
ccl_global char *ray_state, /* Denotes the state of each ray */
int sw, int sh,
ccl_global int *Queue_data, /* Memory for queues */
@@ -121,7 +121,7 @@ __kernel void kernel_ocl_path_trace_SceneIntersect_SPLIT_KERNEL(
#ifdef __KERNEL_DEBUG__
ccl_global DebugData *debug_data = &debugdata_coop[ray_index];
#endif
- ccl_global Intersection *isect = &Intersection_coop[ray_index];
+ Intersection *isect = &Intersection_coop[ray_index];
PathState state = PathState_coop[ray_index];
Ray ray = Ray_coop[ray_index];
diff --git a/intern/cycles/kernel/kernel_ShaderEval.cl b/intern/cycles/kernel/kernel_ShaderEval.cl
index de4b0e3..096cbc1 100644
--- a/intern/cycles/kernel/kernel_ShaderEval.cl
+++ b/intern/cycles/kernel/kernel_ShaderEval.cl
@@ -53,7 +53,7 @@ __kernel void kernel_ocl_path_trace_ShaderEvaluation_SPLIT_KERNEL(
ccl_global uint *rng_coop, /* Required for rbsdf calculation */
ccl_global Ray *Ray_coop, /* Required for setting up shader from ray */
ccl_global PathState *PathState_coop, /* Required for all functions in this kernel */
- ccl_global Intersection *Intersection_coop, /* Required for setting up shader from ray */
+ Intersection *Intersection_coop, /* Required for setting up shader from ray */
ccl_global char *ray_state, /* Denotes the state of each ray */
ccl_global int *Queue_data, /* queue memory */
ccl_global int *Queue_index, /* Tracks the number of elements in each queue */
@@ -81,7 +81,7 @@ __kernel void kernel_ocl_path_trace_ShaderEvaluation_SPLIT_KERNEL(
if(IS_STATE(ray_state, ray_index, RAY_ACTIVE)) {
KernelGlobals *kg = (KernelGlobals *)globals;
ShaderData *sd = (ShaderData *)shader_data;
- ccl_global Intersection *isect = &Intersection_coop[ray_index];
+ Intersection *isect = &Intersection_coop[ray_index];
ccl_global uint *rng = &rng_coop[ray_index];
ccl_global PathState *state = &PathState_coop[ray_index];
Ray ray = Ray_coop[ray_index];
diff --git a/intern/cycles/kernel/kernel_ShadowBlocked.cl b/intern/cycles/kernel/kernel_ShadowBlocked.cl
index b88e131..8521711 100644
--- a/intern/cycles/kernel/kernel_ShadowBlocked.cl
+++ b/intern/cycles/kernel/kernel_ShadowBlocked.cl
@@ -54,8 +54,8 @@ __kernel void kernel_ocl_path_trace_ShadowBlocked_DirectLighting_SPLIT_KERNEL(
ccl_global PathState *PathState_coop, /* Required for shadow blocked */
ccl_global Ray *LightRay_dl_coop, /* Required for direct lighting's shadow blocked */
ccl_global Ray *LightRay_ao_coop, /* Required for AO's shadow blocked */
- ccl_global Intersection *Intersection_coop_AO,
-
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list