[Bf-blender-cvs] [f491c23f1e1] master: Cycles: inline more functions on the GPU
Patrick Mours
noreply at git.blender.org
Mon Aug 26 10:35:25 CEST 2019
Commit: f491c23f1e104998752380b930307e7abc4597b3
Author: Patrick Mours
Date: Thu Aug 22 17:36:54 2019 +0200
Branches: master
https://developer.blender.org/rBf491c23f1e104998752380b930307e7abc4597b3
Cycles: inline more functions on the GPU
This makes little difference for CUDA and OpenCL, but will be helpful
for Optix.
===================================================================
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_path_branched.h
M intern/cycles/kernel/kernel_path_surface.h
M intern/cycles/kernel/kernel_path_volume.h
M intern/cycles/kernel/kernel_volume.h
M intern/cycles/kernel/svm/svm_attribute.h
M intern/cycles/kernel/svm/svm_brick.h
M intern/cycles/kernel/svm/svm_color_util.h
M intern/cycles/kernel/svm/svm_magic.h
M intern/cycles/kernel/svm/svm_musgrave.h
M intern/cycles/kernel/svm/svm_noise.h
M intern/cycles/kernel/svm/svm_wave.h
M intern/cycles/util/util_defines.h
===================================================================
diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h
index 469b81d120b..5075c434b10 100644
--- a/intern/cycles/kernel/kernel_compat_cuda.h
+++ b/intern/cycles/kernel/kernel_compat_cuda.h
@@ -58,6 +58,7 @@ __device__ half __float2half(const float f)
# define ccl_device_forceinline __device__ __forceinline__
#endif
#define ccl_device_noinline __device__ __noinline__
+#define ccl_device_noinline_cpu ccl_device
#define ccl_global
#define ccl_static_constant __constant__
#define ccl_constant const
diff --git a/intern/cycles/kernel/kernel_compat_opencl.h b/intern/cycles/kernel/kernel_compat_opencl.h
index e040ea88d7c..1fe52c51ab0 100644
--- a/intern/cycles/kernel/kernel_compat_opencl.h
+++ b/intern/cycles/kernel/kernel_compat_opencl.h
@@ -35,6 +35,7 @@
#define ccl_device_inline ccl_device
#define ccl_device_forceinline ccl_device
#define ccl_device_noinline ccl_device ccl_noinline
+#define ccl_device_noinline_cpu ccl_device
#define ccl_may_alias
#define ccl_static_constant static __constant
#define ccl_constant __constant
diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h
index be0a2bd2d6b..16d52b0c733 100644
--- a/intern/cycles/kernel/kernel_emission.h
+++ b/intern/cycles/kernel/kernel_emission.h
@@ -17,14 +17,14 @@
CCL_NAMESPACE_BEGIN
/* Direction Emission */
-ccl_device_noinline float3 direct_emissive_eval(KernelGlobals *kg,
- ShaderData *emission_sd,
- LightSample *ls,
- ccl_addr_space PathState *state,
- float3 I,
- differential3 dI,
- float t,
- float time)
+ccl_device_noinline_cpu float3 direct_emissive_eval(KernelGlobals *kg,
+ ShaderData *emission_sd,
+ LightSample *ls,
+ ccl_addr_space PathState *state,
+ float3 I,
+ differential3 dI,
+ float t,
+ float time)
{
/* setup shading at emitter */
float3 eval = make_float3(0.0f, 0.0f, 0.0f);
@@ -98,15 +98,15 @@ ccl_device_noinline float3 direct_emissive_eval(KernelGlobals *kg,
return eval;
}
-ccl_device_noinline bool direct_emission(KernelGlobals *kg,
- ShaderData *sd,
- ShaderData *emission_sd,
- LightSample *ls,
- ccl_addr_space PathState *state,
- Ray *ray,
- BsdfEval *eval,
- bool *is_lamp,
- float rand_terminate)
+ccl_device_noinline_cpu bool direct_emission(KernelGlobals *kg,
+ ShaderData *sd,
+ ShaderData *emission_sd,
+ LightSample *ls,
+ ccl_addr_space PathState *state,
+ Ray *ray,
+ BsdfEval *eval,
+ bool *is_lamp,
+ float rand_terminate)
{
if (ls->pdf == 0.0f)
return false;
@@ -208,7 +208,7 @@ ccl_device_noinline bool direct_emission(KernelGlobals *kg,
/* Indirect Primitive Emission */
-ccl_device_noinline float3 indirect_primitive_emission(
+ccl_device_noinline_cpu float3 indirect_primitive_emission(
KernelGlobals *kg, ShaderData *sd, float t, int path_flag, float bsdf_pdf)
{
/* evaluate emissive closure */
@@ -234,11 +234,11 @@ ccl_device_noinline float3 indirect_primitive_emission(
/* Indirect Lamp Emission */
-ccl_device_noinline bool indirect_lamp_emission(KernelGlobals *kg,
- ShaderData *emission_sd,
- ccl_addr_space PathState *state,
- Ray *ray,
- float3 *emission)
+ccl_device_noinline_cpu bool indirect_lamp_emission(KernelGlobals *kg,
+ ShaderData *emission_sd,
+ ccl_addr_space PathState *state,
+ Ray *ray,
+ float3 *emission)
{
bool hit_lamp = false;
@@ -293,10 +293,10 @@ ccl_device_noinline bool indirect_lamp_emission(KernelGlobals *kg,
/* Indirect Background */
-ccl_device_noinline float3 indirect_background(KernelGlobals *kg,
- ShaderData *emission_sd,
- ccl_addr_space PathState *state,
- ccl_addr_space Ray *ray)
+ccl_device_noinline_cpu float3 indirect_background(KernelGlobals *kg,
+ ShaderData *emission_sd,
+ ccl_addr_space PathState *state,
+ ccl_addr_space Ray *ray)
{
#ifdef __BACKGROUND__
int shader = kernel_data.background.surface_shader;
diff --git a/intern/cycles/kernel/kernel_light.h b/intern/cycles/kernel/kernel_light.h
index 9128bfa9d95..758e91159b6 100644
--- a/intern/cycles/kernel/kernel_light.h
+++ b/intern/cycles/kernel/kernel_light.h
@@ -182,17 +182,7 @@ ccl_device float lamp_light_pdf(KernelGlobals *kg, const float3 Ng, const float3
#ifdef __BACKGROUND_MIS__
-/* TODO(sergey): In theory it should be all fine to use noinline for all
- * devices, but we're so close to the release so better not screw things
- * up for CPU at least.
- */
-# ifdef __KERNEL_GPU__
-ccl_device_noinline
-# else
-ccl_device
-# endif
- float3
- background_map_sample(KernelGlobals *kg, float randu, float randv, float *pdf)
+ccl_device float3 background_map_sample(KernelGlobals *kg, float randu, float randv, float *pdf)
{
/* for the following, the CDF values are actually a pair of floats, with the
* function value as X and the actual CDF as Y. The last entry's function
@@ -274,13 +264,7 @@ ccl_device
/* TODO(sergey): Same as above, after the release we should consider using
* 'noinline' for all devices.
*/
-# ifdef __KERNEL_GPU__
-ccl_device_noinline
-# else
-ccl_device
-# endif
- float
- background_map_pdf(KernelGlobals *kg, float3 direction)
+ccl_device float background_map_pdf(KernelGlobals *kg, float3 direction)
{
float2 uv = direction_to_equirectangular(direction);
int res_x = kernel_data.integrator.pdf_background_res_x;
diff --git a/intern/cycles/kernel/kernel_path_branched.h b/intern/cycles/kernel/kernel_path_branched.h
index 737a7c4aa84..ea6b23e7eb4 100644
--- a/intern/cycles/kernel/kernel_path_branched.h
+++ b/intern/cycles/kernel/kernel_path_branched.h
@@ -198,14 +198,14 @@ ccl_device_forceinline void kernel_branched_path_volume(KernelGlobals *kg,
# endif /* __VOLUME__ */
/* bounce off surface and integrate indirect light */
-ccl_device_noinline void kernel_branched_path_surface_indirect_light(KernelGlobals *kg,
- ShaderData *sd,
- ShaderData *indirect_sd,
- ShaderData *emission_sd,
- float3 throughput,
- float num_samples_adjust,
- PathState *state,
- PathRadiance *L)
+ccl_device_noinline_cpu void kernel_branched_path_surface_indirect_light(KernelGlobals *kg,
+ ShaderData *sd,
+ ShaderData *indirect_sd,
+ ShaderData *emission_sd,
+ float3 throughput,
+ float num_samples_adjust,
+ PathState *state,
+ PathRadiance *L)
{
float sum_sample_weight = 0.0f;
# ifdef __DENOISING_FEATURES__
diff --git a/intern/cycles/kernel/kernel_path_surface.h b/intern/cycles/kernel/kernel_path_surface.h
index a1ab4951565..07444a98d82 100644
--- a/intern/cycles/kernel/kernel_path_surface.h
+++ b/intern/cycles/kernel/kernel_path_surface.h
@@ -20,7 +20,7 @@ CCL_NAMESPACE_BEGIN
defined(__BAKING__)
/* branched path tracing: connect path directly to position on one or more lights and add it to L
*/
-ccl_device_noinline void kernel_branched_path_surface_connect_light(
+ccl_device_noinline_cpu void kernel_branched_path_surface_connect_light(
KernelGlobals *kg,
ShaderData *sd,
ShaderData *emission_sd,
diff --git a/intern/cycles/kernel/kernel_path_volume.h b/intern/cycles/kernel/kernel_path_volume.h
index fea4dfc159d..82975c2ad26 100644
--- a/intern/cycles/kernel/kernel_path_volume.h
+++ b/intern/cycles/kernel/kernel_path_volume.h
@@ -57,18 +57,12 @@ ccl_device_inline void kernel_path_volume_connect_light(KernelG
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list