[Bf-blender-cvs] [eca7e3f8da1] cycles-x: Cycles X: slightly simplify noinline usage
Brecht Van Lommel
noreply at git.blender.org
Wed Jun 2 18:03:45 CEST 2021
Commit: eca7e3f8da1354e6ba2ea517bb09ca549a0dba7f
Author: Brecht Van Lommel
Date: Wed Jun 2 14:05:01 2021 +0200
Branches: cycles-x
https://developer.blender.org/rBeca7e3f8da1354e6ba2ea517bb09ca549a0dba7f
Cycles X: slightly simplify noinline usage
No performance difference found with this. Making this change to help identify
where it's actually needed.
===================================================================
M intern/cycles/kernel/integrator/integrator_shade_background.h
M intern/cycles/kernel/kernel_bake.h
M intern/cycles/kernel/kernel_volume.h
M intern/cycles/kernel/svm/svm.h
M intern/cycles/kernel/svm/svm_ao.h
M intern/cycles/kernel/svm/svm_bevel.h
M intern/cycles/kernel/svm/svm_noise.h
===================================================================
diff --git a/intern/cycles/kernel/integrator/integrator_shade_background.h b/intern/cycles/kernel/integrator/integrator_shade_background.h
index 4b81fdc347f..b3559a1b04c 100644
--- a/intern/cycles/kernel/integrator/integrator_shade_background.h
+++ b/intern/cycles/kernel/integrator/integrator_shade_background.h
@@ -23,8 +23,8 @@
CCL_NAMESPACE_BEGIN
-ccl_device_noinline_cpu float3 integrator_eval_background_shader(
- INTEGRATOR_STATE_ARGS, ccl_global float *ccl_restrict render_buffer)
+ccl_device float3 integrator_eval_background_shader(INTEGRATOR_STATE_ARGS,
+ ccl_global float *ccl_restrict render_buffer)
{
#ifdef __BACKGROUND__
const int shader = kernel_data.background.surface_shader;
diff --git a/intern/cycles/kernel/kernel_bake.h b/intern/cycles/kernel/kernel_bake.h
index 1556b482667..adad35be4c2 100644
--- a/intern/cycles/kernel/kernel_bake.h
+++ b/intern/cycles/kernel/kernel_bake.h
@@ -27,7 +27,7 @@ CCL_NAMESPACE_BEGIN
#if 0
# ifdef __BAKING__
-ccl_device_noinline void compute_light_pass(const KernelGlobals *kg,
+ccl_device void compute_light_pass(const KernelGlobals *kg,
ShaderData *sd,
PathRadiance *L,
uint rng_hash,
diff --git a/intern/cycles/kernel/kernel_volume.h b/intern/cycles/kernel/kernel_volume.h
index 4575d81a367..a2bcfa1aa51 100644
--- a/intern/cycles/kernel/kernel_volume.h
+++ b/intern/cycles/kernel/kernel_volume.h
@@ -289,7 +289,7 @@ ccl_device_inline void kernel_volume_shadow(const KernelGlobals *kg,
}
extern "C" __device__ void __direct_callable__kernel_volume_shadow(
# else
-ccl_device_noinline void kernel_volume_shadow(
+ccl_device void kernel_volume_shadow(
# endif
const KernelGlobals *kg,
ShaderData *shadow_sd,
diff --git a/intern/cycles/kernel/svm/svm.h b/intern/cycles/kernel/svm/svm.h
index 13affa5640b..30e0b78d659 100644
--- a/intern/cycles/kernel/svm/svm.h
+++ b/intern/cycles/kernel/svm/svm.h
@@ -231,10 +231,10 @@ extern "C" __device__ void __direct_callable__svm_eval_nodes(
#endif
template<uint node_feature_mask, ShaderType type>
-ccl_device_noinline void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS,
- ShaderData *sd,
- ccl_global float *buffer,
- int path_flag)
+ccl_device void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS,
+ ShaderData *sd,
+ ccl_global float *buffer,
+ int path_flag)
{
float stack[SVM_STACK_SIZE];
int offset = sd->shader & SHADER_MASK;
diff --git a/intern/cycles/kernel/svm/svm_ao.h b/intern/cycles/kernel/svm/svm_ao.h
index 1cff8d678b3..0e8cc933b9f 100644
--- a/intern/cycles/kernel/svm/svm_ao.h
+++ b/intern/cycles/kernel/svm/svm_ao.h
@@ -20,12 +20,12 @@ CCL_NAMESPACE_BEGIN
#ifdef __SHADER_RAYTRACE__
-ccl_device_noinline float svm_ao(INTEGRATOR_STATE_CONST_ARGS,
- ShaderData *sd,
- float3 N,
- float max_dist,
- int num_samples,
- int flags)
+ccl_device float svm_ao(INTEGRATOR_STATE_CONST_ARGS,
+ ShaderData *sd,
+ float3 N,
+ float max_dist,
+ int num_samples,
+ int flags)
{
if (flags & NODE_AO_GLOBAL_RADIUS) {
max_dist = kernel_data.background.ao_distance;
diff --git a/intern/cycles/kernel/svm/svm_bevel.h b/intern/cycles/kernel/svm/svm_bevel.h
index 9ba9182e6c8..2a51e09ee5a 100644
--- a/intern/cycles/kernel/svm/svm_bevel.h
+++ b/intern/cycles/kernel/svm/svm_bevel.h
@@ -28,10 +28,10 @@ CCL_NAMESPACE_BEGIN
* http://library.imageworks.com/pdfs/imageworks-library-BSSRDF-sampling.pdf
*/
-ccl_device_noinline float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
- ShaderData *sd,
- float radius,
- int num_samples)
+ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
+ ShaderData *sd,
+ float radius,
+ int num_samples)
{
/* Early out if no sampling needed. */
if (radius <= 0.0f || num_samples < 1 || sd->object == OBJECT_NONE) {
diff --git a/intern/cycles/kernel/svm/svm_noise.h b/intern/cycles/kernel/svm/svm_noise.h
index 94d8bfde555..ecb4df6afdf 100644
--- a/intern/cycles/kernel/svm/svm_noise.h
+++ b/intern/cycles/kernel/svm/svm_noise.h
@@ -330,7 +330,7 @@ ccl_device_inline ssef grad(const ssei &hash, const ssef &x, const ssef &y)
* |__________________________|
*
*/
-ccl_device_noinline float perlin_2d(float x, float y)
+ccl_device_noinline_cpu float perlin_2d(float x, float y)
{
ssei XY;
ssef fxy = floorfrac(ssef(x, y, 0.0f, 0.0f), &XY);
@@ -447,7 +447,7 @@ ccl_device_inline ssef quad_mix(ssef p, ssef q, ssef r, ssef s, ssef f)
* v7 (1, 1, 1)
*
*/
-ccl_device_noinline float perlin_3d(float x, float y, float z)
+ccl_device_noinline_cpu float perlin_3d(float x, float y, float z)
{
ssei XYZ;
ssef fxyz = floorfrac(ssef(x, y, z, 0.0f), &XYZ);
@@ -501,7 +501,7 @@ ccl_device_noinline float perlin_3d(float x, float y, float z)
* v15 (1, 1, 1, 1)
*
*/
-ccl_device_noinline float perlin_4d(float x, float y, float z, float w)
+ccl_device_noinline_cpu float perlin_4d(float x, float y, float z, float w)
{
ssei XYZW;
ssef fxyzw = floorfrac(ssef(x, y, z, w), &XYZW);
@@ -585,7 +585,7 @@ ccl_device_inline ssef quad_mix(avxf p, avxf q, ssef f)
* |__________________________|
*
*/
-ccl_device_noinline float perlin_3d(float x, float y, float z)
+ccl_device_noinline_cpu float perlin_3d(float x, float y, float z)
{
ssei XYZ;
ssef fxyz = floorfrac(ssef(x, y, z, 0.0f), &XYZ);
@@ -637,7 +637,7 @@ ccl_device_noinline float perlin_3d(float x, float y, float z)
* v15 (1, 1, 1, 1)
*
*/
-ccl_device_noinline float perlin_4d(float x, float y, float z, float w)
+ccl_device_noinline_cpu float perlin_4d(float x, float y, float z, float w)
{
ssei XYZW;
ssef fxyzw = floorfrac(ssef(x, y, z, w), &XYZW);
More information about the Bf-blender-cvs
mailing list