[Bf-blender-cvs] [570939529e0] cycles_refactor: Cycles: Added kernel_struct_fetch() macro
Stefan Werner
noreply at git.blender.org
Thu Feb 1 13:59:30 CET 2018
Commit: 570939529e090bdcc7075d6028a83bbed1dac0b5
Author: Stefan Werner
Date: Wed Dec 6 09:32:33 2017 +0100
Branches: cycles_refactor
https://developer.blender.org/rB570939529e090bdcc7075d6028a83bbed1dac0b5
Cycles: Added kernel_struct_fetch() macro
===================================================================
M intern/cycles/kernel/bvh/bvh_shadow_all.h
M intern/cycles/kernel/bvh/qbvh_shadow_all.h
M intern/cycles/kernel/geom/geom_attribute.h
M intern/cycles/kernel/geom/geom_object.h
M intern/cycles/kernel/kernel_compat_cpu.h
M intern/cycles/kernel/kernel_compat_cuda.h
M intern/cycles/kernel/kernel_compat_opencl.h
M intern/cycles/kernel/kernel_emission.h
M intern/cycles/kernel/kernel_light.h
M intern/cycles/kernel/kernel_shader.h
M intern/cycles/kernel/kernel_volume.h
===================================================================
diff --git a/intern/cycles/kernel/bvh/bvh_shadow_all.h b/intern/cycles/kernel/bvh/bvh_shadow_all.h
index 86cab44eb44..4cdb6934a9a 100644
--- a/intern/cycles/kernel/bvh/bvh_shadow_all.h
+++ b/intern/cycles/kernel/bvh/bvh_shadow_all.h
@@ -276,7 +276,7 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
shader = __float_as_int(str.z);
}
#endif
- int flag = kernel_tex_fetch(__shader_flag, (shader & SHADER_MASK)).flags;
+ int flag = kernel_struct_fetch(__shader_flag, flags, (shader & SHADER_MASK));
/* if no transparent shadows, all light is blocked */
if(!(flag & SD_HAS_TRANSPARENT_SHADOW)) {
diff --git a/intern/cycles/kernel/bvh/qbvh_shadow_all.h b/intern/cycles/kernel/bvh/qbvh_shadow_all.h
index 8c62ba81692..db48408dfb7 100644
--- a/intern/cycles/kernel/bvh/qbvh_shadow_all.h
+++ b/intern/cycles/kernel/bvh/qbvh_shadow_all.h
@@ -358,7 +358,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
shader = __float_as_int(str.z);
}
#endif
- int flag = kernel_tex_fetch(__shader_flag, (shader & SHADER_MASK)).flags;
+ int flag = kernel_struct_fetch(__shader_flag, flags, (shader & SHADER_MASK));
/* if no transparent shadows, all light is blocked */
if(!(flag & SD_HAS_TRANSPARENT_SHADOW)) {
diff --git a/intern/cycles/kernel/geom/geom_attribute.h b/intern/cycles/kernel/geom/geom_attribute.h
index 6e2ee3b62d7..61dcd44a1f9 100644
--- a/intern/cycles/kernel/geom/geom_attribute.h
+++ b/intern/cycles/kernel/geom/geom_attribute.h
@@ -53,7 +53,7 @@ ccl_device_inline AttributeDescriptor attribute_not_found()
ccl_device_inline uint object_attribute_map_offset(KernelGlobals *kg, int object)
{
- return kernel_tex_fetch(__objects, object).attribute_map_offset;
+ return kernel_struct_fetch(__objects, attribute_map_offset, object);
}
ccl_device_inline AttributeDescriptor find_attribute(KernelGlobals *kg, const ShaderData *sd, uint id)
diff --git a/intern/cycles/kernel/geom/geom_object.h b/intern/cycles/kernel/geom/geom_object.h
index 714a2176261..edf2ee81d37 100644
--- a/intern/cycles/kernel/geom/geom_object.h
+++ b/intern/cycles/kernel/geom/geom_object.h
@@ -236,7 +236,7 @@ ccl_device_inline float3 object_location(KernelGlobals *kg, const ShaderData *sd
ccl_device_inline float object_surface_area(KernelGlobals *kg, int object)
{
- return kernel_tex_fetch(__objects, object).surface_area;
+ return kernel_struct_fetch(__objects, surface_area, object);
}
/* Pass ID number of object */
@@ -246,7 +246,7 @@ ccl_device_inline float object_pass_id(KernelGlobals *kg, int object)
if(object == OBJECT_NONE)
return 0.0f;
- return kernel_tex_fetch(__objects, object).pass_id;
+ return kernel_struct_fetch(__objects, pass_id, object);
}
/* Per lamp random number for shader variation */
@@ -256,7 +256,7 @@ ccl_device_inline float lamp_random_number(KernelGlobals *kg, int lamp)
if(lamp == LAMP_NONE)
return 0.0f;
- return kernel_tex_fetch(__light_data, lamp).random;
+ return kernel_struct_fetch(__light_data, random, lamp);
}
/* Per object random number for shader variation */
@@ -266,7 +266,7 @@ ccl_device_inline float object_random_number(KernelGlobals *kg, int object)
if(object == OBJECT_NONE)
return 0.0f;
- return kernel_tex_fetch(__objects, object).random_number;
+ return kernel_struct_fetch(__objects, random_number, object);
}
/* Particle ID from which this object was generated */
@@ -276,7 +276,7 @@ ccl_device_inline int object_particle_id(KernelGlobals *kg, int object)
if(object == OBJECT_NONE)
return 0;
- return kernel_tex_fetch(__objects, object).particle_index;
+ return kernel_struct_fetch(__objects, particle_index, object);
}
/* Generated texture coordinate on surface from where object was instanced */
@@ -306,13 +306,13 @@ ccl_device_inline float3 object_dupli_uv(KernelGlobals *kg, int object)
ccl_device_inline void object_motion_info(KernelGlobals *kg, int object, int *numsteps, int *numverts, int *numkeys)
{
if(numkeys) {
- *numkeys = kernel_tex_fetch(__objects, object).numkeys;
+ *numkeys = kernel_struct_fetch(__objects, numkeys, object);
}
if(numsteps)
- *numsteps = kernel_tex_fetch(__objects, object).numsteps;
+ *numsteps = kernel_struct_fetch(__objects, numsteps, object);
if(numverts)
- *numverts = kernel_tex_fetch(__objects, object).numverts;
+ *numverts = kernel_struct_fetch(__objects, numverts, object);
}
/* Offset to an objects patch map */
@@ -322,56 +322,56 @@ ccl_device_inline uint object_patch_map_offset(KernelGlobals *kg, int object)
if(object == OBJECT_NONE)
return 0;
- return kernel_tex_fetch(__objects, object).patch_map_offset;
+ return kernel_struct_fetch(__objects, patch_map_offset, object);
}
/* Pass ID for shader */
ccl_device int shader_pass_id(KernelGlobals *kg, const ShaderData *sd)
{
- return kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)).pass_id;
+ return kernel_struct_fetch(__shader_flag, pass_id, (sd->shader & SHADER_MASK));
}
/* Particle data from which object was instanced */
ccl_device_inline float particle_index(KernelGlobals *kg, int particle)
{
- return kernel_tex_fetch(__particles, particle).index;
+ return kernel_struct_fetch(__particles, index, particle);
}
ccl_device float particle_age(KernelGlobals *kg, int particle)
{
- return kernel_tex_fetch(__particles, particle).age;
+ return kernel_struct_fetch(__particles, age, particle);
}
ccl_device float particle_lifetime(KernelGlobals *kg, int particle)
{
- return kernel_tex_fetch(__particles, particle).lifetime;
+ return kernel_struct_fetch(__particles, lifetime, particle);
}
ccl_device float particle_size(KernelGlobals *kg, int particle)
{
- return kernel_tex_fetch(__particles, particle).size;
+ return kernel_struct_fetch(__particles, size, particle);
}
ccl_device float4 particle_rotation(KernelGlobals *kg, int particle)
{
- return kernel_tex_fetch(__particles, particle).rotation;
+ return kernel_struct_fetch(__particles, rotation, particle);
}
ccl_device float3 particle_location(KernelGlobals *kg, int particle)
{
- return float4_to_float3(kernel_tex_fetch(__particles, particle).location);
+ return float4_to_float3(kernel_struct_fetch(__particles, location, particle));
}
ccl_device float3 particle_velocity(KernelGlobals *kg, int particle)
{
- return float4_to_float3(kernel_tex_fetch(__particles, particle).velocity);
+ return float4_to_float3(kernel_struct_fetch(__particles, velocity, particle));
}
ccl_device float3 particle_angular_velocity(KernelGlobals *kg, int particle)
{
- return float4_to_float3(kernel_tex_fetch(__particles, particle).angular_velocity);
+ return float4_to_float3(kernel_struct_fetch(__particles, angular_velocity, particle));
}
/* Object intersection in BVH */
diff --git a/intern/cycles/kernel/kernel_compat_cpu.h b/intern/cycles/kernel/kernel_compat_cpu.h
index 6f63c8f77a2..7aa4c9fbe4a 100644
--- a/intern/cycles/kernel/kernel_compat_cpu.h
+++ b/intern/cycles/kernel/kernel_compat_cpu.h
@@ -119,6 +119,7 @@ template<typename T> struct texture {
#define kernel_tex_fetch_ssef(tex, index) (kg->tex.fetch_ssef(index))
#define kernel_tex_fetch_ssei(tex, index) (kg->tex.fetch_ssei(index))
#define kernel_tex_lookup(tex, t, offset, size) (kg->tex.lookup(t, offset, size))
+#define kernel_struct_fetch(tex, member, index) (kg->tex.fetch(index).member)
#define kernel_data (kg->__data)
diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h
index fa512f80e41..170453b0b2e 100644
--- a/intern/cycles/kernel/kernel_compat_cuda.h
+++ b/intern/cycles/kernel/kernel_compat_cuda.h
@@ -129,6 +129,11 @@ ccl_device_inline uint ccl_num_groups(uint d)
/* Use arrays for regular data. This is a little slower than textures on Fermi,
* but allows for cleaner code and we will stop supporting Fermi soon. */
#define kernel_tex_fetch(t, index) t[(index)]
+#if __CUDA_ARCH__ < 350
+# define kernel_struct_fetch(t, member, index) (kernel_tex_fetch(t, index).member)
+#else
+# define kernel_struct_fetch(t, member, index) __ldg(&kernel_tex_fetch(t, index).member)
+#endif
/* On Kepler (6xx) and above, we use Bindless Textures for images.
* On Fermi cards (4xx and 5xx), we have to use regular textures. */
diff --git a/intern/cycles/kernel/kernel_compat_opencl.h b/intern/cycles/kernel/kernel_compat_opencl.h
index b02e3bc576d..eb5cd53fce5 100644
--- a/intern/cycles/kernel/kernel_compat_opencl.h
+++ b/intern/cycles/kernel/kernel_compat_opencl.h
@@ -145,6 +145,7 @@
/* data lookup defines */
#define kernel_data (*kg->data)
#define kernel_tex_fetch(tex, index) ((const ccl_global tex##_t*)(kg->buffers[kg->tex.cl_buffer] + kg->tex.data))[(index)]
+#define kernel_struct_fetch(t, member, index) (kernel_tex_fetch(t, index).member)
/* define NULL */
#define NULL 0
diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h
index 651d97b05f8..437d53c254d 100644
--- a/intern/cycles/kernel/kernel_emission.h
+++ b/intern/cycles/kernel/kernel_emission.h
@@ -29,7 +29,7 @@ ccl_device_noinline float3 direct_emissive_eval(KernelGlobals *kg,
/* setup shading at emitter */
float3 eval;
- int shader_flag = kernel_tex_fetch(__shader_flag, (ls->shader & SHADER_MASK)).flags;
+ int shader_flag = kernel_struct_fetch(__shader_flag, flags, (ls->shader & SHADER_MASK));
#ifdef __BACKGROUND_MIS__
if(ls->type == LIGHT_BACKGROUND) {
diff --git a/intern/cycles/kernel/kernel_light.h b/intern/cycles/kernel/kernel_light.h
index 43ebccceaca..bb62af9621e 100644
--- a/intern/cycles/kernel/kernel_light.h
+++ b/intern/cycles/kernel/kernel_light.h
@@ -522,9 +522,9 @@ ccl_device_inline bool lamp_light_sample(KernelGlobals *kg,
float3 P,
LightSample *ls)
{
- LightType type = (LightType)kernel_tex_fetch(__light_data, lamp).type;
+ LightType type = (LightType)kernel_struct_fetch(__light_data, type, lamp);
ls->type = type;
- ls->shader = kernel_tex_fetch(__light_data, lamp).shader_id;
+ ls->shader = kernel_struct_fetch(__light_data, shader_id, lamp);
ls->object = PRIM_NONE;
ls->prim = PRIM_NONE;
ls->lamp =
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list