Cycles: Kernel address space changes for MSL

This is the first of a sequence of changes to support compiling Cycles kernels as MSL (Metal Shading Language) in preparation for a Metal GPU device implementation.

MSL requires that all pointer types be declared with explicit address space attributes (device, thread, etc...). There is already precedent for this with Cycles' address space macros (ccl_global, ccl_private, etc...), therefore the first step of MSL-enablement is to apply these consistently. Line-for-line this represents the largest change required to enable MSL. Applying this change first will simplify future patches as well as offering the emergent benefit of enhanced descriptiveness.

The vast majority of deltas in this patch fall into one of two cases:

- Ensuring ccl_private is specified for thread-local pointer types
- Ensuring ccl_global is specified for device-wide pointer types

Additionally, the ccl_addr_space qualifier can be removed. Prior to Cycles X, ccl_addr_space was used as a context-dependent address space qualifier, but now it is either redundant (e.g. in struct typedefs), or can be replaced by ccl_global in the case of pointer types. Associated function variants (e.g. lcg_step_float_addrspace) are also redundant.

In cases where address space qualifiers are chained with "const", this patch places the address space qualifier first. The rationale for this is that the choice of address space is likely to have the greater impact on runtime performance and overall architecture.

The final part of this patch is the addition of a metal/compat.h header. This is partially complete and will be extended in future patches, paving the way for the full Metal implementation.

Ref T92212

Reviewed By: brecht

Maniphest Tasks: T92212

Differential Revision: https://developer.blender.org/D12864

48 files changed, 508 insertions, 340 deletions

diff --git a/intern/cycles/kernel/svm/svm.h b/intern/cycles/kernel/svm/svm.h
index ad609b15f86..871e370123e 100644
--- a/intern/cycles/kernel/svm/svm.h
+++ b/intern/cycles/kernel/svm/svm.h
@@ -44,56 +44,56 @@ CCL_NAMESPACE_BEGIN
 
 /* Stack */
 
-ccl_device_inline float3 stack_load_float3(float *stack, uint a)
+ccl_device_inline float3 stack_load_float3(ccl_private float *stack, uint a)
 {
   kernel_assert(a + 2 < SVM_STACK_SIZE);
 
-  float *stack_a = stack + a;
+  ccl_private float *stack_a = stack + a;
   return make_float3(stack_a[0], stack_a[1], stack_a[2]);
 }
 
-ccl_device_inline void stack_store_float3(float *stack, uint a, float3 f)
+ccl_device_inline void stack_store_float3(ccl_private float *stack, uint a, float3 f)
 {
   kernel_assert(a + 2 < SVM_STACK_SIZE);
 
-  float *stack_a = stack + a;
+  ccl_private float *stack_a = stack + a;
   stack_a[0] = f.x;
   stack_a[1] = f.y;
   stack_a[2] = f.z;
 }
 
-ccl_device_inline float stack_load_float(float *stack, uint a)
+ccl_device_inline float stack_load_float(ccl_private float *stack, uint a)
 {
   kernel_assert(a < SVM_STACK_SIZE);
 
   return stack[a];
 }
 
-ccl_device_inline float stack_load_float_default(float *stack, uint a, uint value)
+ccl_device_inline float stack_load_float_default(ccl_private float *stack, uint a, uint value)
 {
   return (a == (uint)SVM_STACK_INVALID) ? __uint_as_float(value) : stack_load_float(stack, a);
 }
 
-ccl_device_inline void stack_store_float(float *stack, uint a, float f)
+ccl_device_inline void stack_store_float(ccl_private float *stack, uint a, float f)
 {
   kernel_assert(a < SVM_STACK_SIZE);
 
   stack[a] = f;
 }
 
-ccl_device_inline int stack_load_int(float *stack, uint a)
+ccl_device_inline int stack_load_int(ccl_private float *stack, uint a)
 {
   kernel_assert(a < SVM_STACK_SIZE);
 
   return __float_as_int(stack[a]);
 }
 
-ccl_device_inline int stack_load_int_default(float *stack, uint a, uint value)
+ccl_device_inline int stack_load_int_default(ccl_private float *stack, uint a, uint value)
 {
   return (a == (uint)SVM_STACK_INVALID) ? (int)value : stack_load_int(stack, a);
 }
 
-ccl_device_inline void stack_store_int(float *stack, uint a, int i)
+ccl_device_inline void stack_store_int(ccl_private float *stack, uint a, int i)
 {
   kernel_assert(a < SVM_STACK_SIZE);
 
@@ -107,14 +107,15 @@ ccl_device_inline bool stack_valid(uint a)
 
 /* Reading Nodes */
 
-ccl_device_inline uint4 read_node(const KernelGlobals *kg, int *offset)
+ccl_device_inline uint4 read_node(ccl_global const KernelGlobals *kg, ccl_private int *offset)
 {
   uint4 node = kernel_tex_fetch(__svm_nodes, *offset);
   (*offset)++;
   return node;
 }
 
-ccl_device_inline float4 read_node_float(const KernelGlobals *kg, int *offset)
+ccl_device_inline float4 read_node_float(ccl_global const KernelGlobals *kg,
+                                         ccl_private int *offset)
 {
   uint4 node = kernel_tex_fetch(__svm_nodes, *offset);
   float4 f = make_float4(__uint_as_float(node.x),
@@ -125,7 +126,7 @@ ccl_device_inline float4 read_node_float(const KernelGlobals *kg, int *offset)
   return f;
 }
 
-ccl_device_inline float4 fetch_node_float(const KernelGlobals *kg, int offset)
+ccl_device_inline float4 fetch_node_float(ccl_global const KernelGlobals *kg, int offset)
 {
   uint4 node = kernel_tex_fetch(__svm_nodes, offset);
   return make_float4(__uint_as_float(node.x),
@@ -134,20 +135,26 @@ ccl_device_inline float4 fetch_node_float(const KernelGlobals *kg, int offset)
                      __uint_as_float(node.w));
 }
 
-ccl_device_forceinline void svm_unpack_node_uchar2(uint i, uint *x, uint *y)
+ccl_device_forceinline void svm_unpack_node_uchar2(uint i,
+                                                   ccl_private uint *x,
+                                                   ccl_private uint *y)
 {
   *x = (i & 0xFF);
   *y = ((i >> 8) & 0xFF);
 }
 
-ccl_device_forceinline void svm_unpack_node_uchar3(uint i, uint *x, uint *y, uint *z)
+ccl_device_forceinline void svm_unpack_node_uchar3(uint i,
+                                                   ccl_private uint *x,
+                                                   ccl_private uint *y,
+                                                   ccl_private uint *z)
 {
   *x = (i & 0xFF);
   *y = ((i >> 8) & 0xFF);
   *z = ((i >> 16) & 0xFF);
 }
 
-ccl_device_forceinline void svm_unpack_node_uchar4(uint i, uint *x, uint *y, uint *z, uint *w)
+ccl_device_forceinline void svm_unpack_node_uchar4(
+    uint i, ccl_private uint *x, ccl_private uint *y, ccl_private uint *z, ccl_private uint *w)
 {
   *x = (i & 0xFF);
   *y = ((i >> 8) & 0xFF);
diff --git a/intern/cycles/kernel/svm/svm_ao.h b/intern/cycles/kernel/svm/svm_ao.h
index 34ac2cb8fbf..092f3817fd8 100644
--- a/intern/cycles/kernel/svm/svm_ao.h
+++ b/intern/cycles/kernel/svm/svm_ao.h
@@ -25,7 +25,7 @@ extern "C" __device__ float __direct_callable__svm_node_ao(INTEGRATOR_STATE_CONS
 #  else
 ccl_device float svm_ao(INTEGRATOR_STATE_CONST_ARGS,
 #  endif
-                                                           ShaderData *sd,
+                                                           ccl_private ShaderData *sd,
                                                            float3 N,
                                                            float max_dist,
                                                            int num_samples,
@@ -96,7 +96,10 @@ ccl_device_inline
 ccl_device_noinline
 #  endif
     void
-    svm_node_ao(INTEGRATOR_STATE_CONST_ARGS, ShaderData *sd, float *stack, uint4 node)
+    svm_node_ao(INTEGRATOR_STATE_CONST_ARGS,
+                ccl_private ShaderData *sd,
+                ccl_private float *stack,
+                uint4 node)
 {
   uint flags, dist_offset, normal_offset, out_ao_offset;
   svm_unpack_node_uchar4(node.y, &flags, &dist_offset, &normal_offset, &out_ao_offset);
diff --git a/intern/cycles/kernel/svm/svm_aov.h b/intern/cycles/kernel/svm/svm_aov.h
index 26dec9717b3..640bec87ac9 100644
--- a/intern/cycles/kernel/svm/svm_aov.h
+++ b/intern/cycles/kernel/svm/svm_aov.h
@@ -26,8 +26,8 @@ ccl_device_inline bool svm_node_aov_check(const int path_flag, ccl_global float
 }
 
 ccl_device void svm_node_aov_color(INTEGRATOR_STATE_CONST_ARGS,
-                                   ShaderData *sd,
-                                   float *stack,
+                                   ccl_private ShaderData *sd,
+                                   ccl_private float *stack,
                                    uint4 node,
                                    ccl_global float *render_buffer)
 {
@@ -44,8 +44,8 @@ ccl_device void svm_node_aov_color(INTEGRATOR_STATE_CONST_ARGS,
 }
 
 ccl_device void svm_node_aov_value(INTEGRATOR_STATE_CONST_ARGS,
-                                   ShaderData *sd,
-                                   float *stack,
+                                   ccl_private ShaderData *sd,
+                                   ccl_private float *stack,
                                    uint4 node,
                                    ccl_global float *render_buffer)
 {
diff --git a/intern/cycles/kernel/svm/svm_attribute.h b/intern/cycles/kernel/svm/svm_attribute.h
index 5f94b20af73..9fd401ba1c3 100644
--- a/intern/cycles/kernel/svm/svm_attribute.h
+++ b/intern/cycles/kernel/svm/svm_attribute.h
@@ -18,11 +18,11 @@ CCL_NAMESPACE_BEGIN
 
 /* Attribute Node */
 
-ccl_device AttributeDescriptor svm_node_attr_init(const KernelGlobals *kg,
-                                                  ShaderData *sd,
+ccl_device AttributeDescriptor svm_node_attr_init(ccl_global const KernelGlobals *kg,
+                                                  ccl_private ShaderData *sd,
                                                   uint4 node,
-                                                  NodeAttributeOutputType *type,
-                                                  uint *out_offset)
+                                                  ccl_private NodeAttributeOutputType *type,
+                                                  ccl_private uint *out_offset)
 {
   *out_offset = node.z;
   *type = (NodeAttributeOutputType)node.w;
@@ -48,9 +48,9 @@ ccl_device AttributeDescriptor svm_node_attr_init(const KernelGlobals *kg,
 }
 
 template<uint node_feature_mask>
-ccl_device_noinline void svm_node_attr(const KernelGlobals *kg,
-                                       ShaderData *sd,
-                                       float *stack,
+ccl_device_noinline void svm_node_attr(ccl_global const KernelGlobals *kg,
+                                       ccl_private ShaderData *sd,
+                                       ccl_private float *stack,
                                        uint4 node)
 {
   NodeAttributeOutputType type = NODE_ATTR_OUTPUT_FLOAT;
@@ -148,9 +148,9 @@ ccl_device_noinline void svm_node_attr(const KernelGlobals *kg,
   }
 }
 
-ccl_device_noinline void svm_node_attr_bump_dx(const KernelGlobals *kg,
-                                               ShaderData *sd,
-                                               float *stack,
+ccl_device_noinline void svm_node_attr_bump_dx(ccl_global const KernelGlobals *kg,
+                                               ccl_private ShaderData *sd,
+                                               ccl_private float *stack,
                                                uint4 node)
 {
   NodeAttributeOutputType type = NODE_ATTR_OUTPUT_FLOAT;
@@ -244,9 +244,9 @@ ccl_device_noinline void svm_node_attr_bump_dx(const KernelGlobals *kg,
   }
 }
 
-ccl_device_noinline void svm_node_attr_bump_dy(const KernelGlobals *kg,
-                                               ShaderData *sd,
-                                               float *stack,
+ccl_device_noinline void svm_node_attr_bump_dy(ccl_global const KernelGlobals *kg,
+                                               ccl_private ShaderData *sd,
+                                               ccl_private float *stack,
                                                uint4 node)
 {
   NodeAttributeOutputType type = NODE_ATTR_OUTPUT_FLOAT;
diff --git a/intern/cycles/kernel/svm/svm_bevel.h b/intern/cycles/kernel/svm/svm_bevel.h
index 60302b8e3d7..a76584e6bc8 100644
--- a/intern/cycles/kernel/svm/svm_bevel.h
+++ b/intern/cycles/kernel/svm/svm_bevel.h
@@ -77,7 +77,10 @@ ccl_device_forceinline float svm_bevel_cubic_quintic_root_find(float xi)
   return x;
 }
 
-ccl_device void svm_bevel_cubic_sample(const float radius, float xi, float *r, float *h)
+ccl_device void svm_bevel_cubic_sample(const float radius,
+                                       float xi,
+                                       ccl_private float *r,
+                                       ccl_private float *h)
 {
   float Rm = radius;
   float r_ = svm_bevel_cubic_quintic_root_find(xi);
@@ -100,7 +103,7 @@ extern "C" __device__ float3 __direct_callable__svm_node_bevel(INTEGRATOR_STATE_
 #  else
 ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
 #  endif
-                                                               ShaderData *sd,
+                                                               ccl_private ShaderData *sd,
                                                                float radius,
                                                                int num_samples)
 {
@@ -284,7 +287,10 @@ ccl_device_inline
 ccl_device_noinline
 #  endif
     void
-    svm_node_bevel(INTEGRATOR_STATE_CONST_ARGS, ShaderData *sd, float *stack, uint4 node)
+    svm_node_bevel(INTEGRATOR_STATE_CONST_ARGS,
+                   ccl_private ShaderData *sd,
+                   ccl_private float *stack,
+                   uint4 node)
 {
   uint num_samples, radius_offset, normal_offset, out_offset;
   svm_unpack_node_uchar4(node.y, &num_samples, &radius_offset, &normal_offset, &out_offset);
diff --git a/intern/cycles/kernel/svm/svm_blackbody.h b/intern/cycles/kernel/svm/svm_blackbody.h
index 96b3703b954..521afb42adc 100644
--- a/intern/cycles/kernel/svm/svm_blackbody.h
+++ b/intern/cycles/kernel/svm/svm_blackbody.h
@@ -34,9 +34,9 @@ CCL_NAMESPACE_BEGIN
 
 /* Blackbody Node */
 
-ccl_device_noinline void svm_node_blackbody(const KernelGlobals *kg,
-                                            ShaderData *sd,
-                                            float *stack,
+ccl_device_noinline void svm_node_blackbody(ccl_global const KernelGlobals *kg,
+                                            ccl_private ShaderData *sd,
+                                            ccl_private float *stack,
                                             uint temperature_offset,
                                             uint col_offset)
 {
diff --git a/intern/cycles/kernel/svm/svm_brick.h b/intern/cycles/kernel/svm/svm_brick.h
index dca1b220dd5..29a8350f1c1 100644
--- a/intern/cycles/kernel/svm/svm_brick.h
+++ b/intern/cycles/kernel/svm/svm_brick.h
@@ -72,8 +72,11 @@ ccl_device_noinline_cpu float2 svm_brick(float3 p,
   return make_float2(tint, mortar);
 }
 
-ccl_device_noinline int svm_node_tex_brick(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_tex_brick(ccl_global const KernelGlobals *kg,
+                                           ccl_private ShaderData *sd,
+                                           ccl_private float *stack,
+                                           uint4 node,
+                                           int offset)
 {
   uint4 node2 = read_node(kg, &offset);
   uint4 node3 = read_node(kg, &offset);
diff --git a/intern/cycles/kernel/svm/svm_brightness.h b/intern/cycles/kernel/svm/svm_brightness.h
index 2ed812acd71..0a44ffe6359 100644
--- a/intern/cycles/kernel/svm/svm_brightness.h
+++ b/intern/cycles/kernel/svm/svm_brightness.h
@@ -17,7 +17,7 @@
 CCL_NAMESPACE_BEGIN
 
 ccl_device_noinline void svm_node_brightness(
-    ShaderData *sd, float *stack, uint in_color, uint out_color, uint node)
+    ccl_private ShaderData *sd, ccl_private float *stack, uint in_color, uint out_color, uint node)
 {
   uint bright_offset, contrast_offset;
   float3 color = stack_load_float3(stack, in_color);
diff --git a/intern/cycles/kernel/svm/svm_bump.h b/intern/cycles/kernel/svm/svm_bump.h
index 8672839dbab..70935c730f4 100644
--- a/intern/cycles/kernel/svm/svm_bump.h
+++ b/intern/cycles/kernel/svm/svm_bump.h
@@ -18,9 +18,9 @@ CCL_NAMESPACE_BEGIN
 
 /* Bump Eval Nodes */
 
-ccl_device_noinline void svm_node_enter_bump_eval(const KernelGlobals *kg,
-                                                  ShaderData *sd,
-                                                  float *stack,
+ccl_device_noinline void svm_node_enter_bump_eval(ccl_global const KernelGlobals *kg,
+                                                  ccl_private ShaderData *sd,
+                                                  ccl_private float *stack,
                                                   uint offset)
 {
   /* save state */
@@ -45,9 +45,9 @@ ccl_device_noinline void svm_node_enter_bump_eval(const KernelGlobals *kg,
   }
 }
 
-ccl_device_noinline void svm_node_leave_bump_eval(const KernelGlobals *kg,
-                                                  ShaderData *sd,
-                                                  float *stack,
+ccl_device_noinline void svm_node_leave_bump_eval(ccl_global const KernelGlobals *kg,
+                                                  ccl_private ShaderData *sd,
+                                                  ccl_private float *stack,
                                                   uint offset)
 {
   /* restore state */
diff --git a/intern/cycles/kernel/svm/svm_camera.h b/intern/cycles/kernel/svm/svm_camera.h
index 40c0edcdad0..2b786757af8 100644
--- a/intern/cycles/kernel/svm/svm_camera.h
+++ b/intern/cycles/kernel/svm/svm_camera.h
@@ -16,9 +16,9 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device_noinline void svm_node_camera(const KernelGlobals *kg,
-                                         ShaderData *sd,
-                                         float *stack,
+ccl_device_noinline void svm_node_camera(ccl_global const KernelGlobals *kg,
+                                         ccl_private ShaderData *sd,
+                                         ccl_private float *stack,
                                          uint out_vector,
                                          uint out_zdepth,
                                          uint out_distance)
diff --git a/intern/cycles/kernel/svm/svm_checker.h b/intern/cycles/kernel/svm/svm_checker.h
index a9919c9ddc9..e22367f4f59 100644
--- a/intern/cycles/kernel/svm/svm_checker.h
+++ b/intern/cycles/kernel/svm/svm_checker.h
@@ -32,9 +32,9 @@ ccl_device float svm_checker(float3 p)
   return ((xi % 2 == yi % 2) == (zi % 2)) ? 1.0f : 0.0f;
 }
 
-ccl_device_noinline void svm_node_tex_checker(const KernelGlobals *kg,
-                                              ShaderData *sd,
-                                              float *stack,
+ccl_device_noinline void svm_node_tex_checker(ccl_global const KernelGlobals *kg,
+                                              ccl_private ShaderData *sd,
+                                              ccl_private float *stack,
                                               uint4 node)
 {
   uint co_offset, color1_offset, color2_offset, scale_offset;
diff --git a/intern/cycles/kernel/svm/svm_clamp.h b/intern/cycles/kernel/svm/svm_clamp.h
index 656bd31c085..cb5224aebb2 100644
--- a/intern/cycles/kernel/svm/svm_clamp.h
+++ b/intern/cycles/kernel/svm/svm_clamp.h
@@ -18,9 +18,9 @@ CCL_NAMESPACE_BEGIN
 
 /* Clamp Node */
 
-ccl_device_noinline int svm_node_clamp(const KernelGlobals *kg,
-                                       ShaderData *sd,
-                                       float *stack,
+ccl_device_noinline int svm_node_clamp(ccl_global const KernelGlobals *kg,
+                                       ccl_private ShaderData *sd,
+                                       ccl_private float *stack,
                                        uint value_stack_offset,
                                        uint parameters_stack_offsets,
                                        uint result_stack_offset,
diff --git a/intern/cycles/kernel/svm/svm_closure.h b/intern/cycles/kernel/svm/svm_closure.h
index e55f76a4400..87be73bb2cc 100644
--- a/intern/cycles/kernel/svm/svm_closure.h
+++ b/intern/cycles/kernel/svm/svm_closure.h
@@ -18,8 +18,12 @@ CCL_NAMESPACE_BEGIN
 
 /* Closure Nodes */
 
-ccl_device void svm_node_glass_setup(
-    ShaderData *sd, MicrofacetBsdf *bsdf, int type, float eta, float roughness, bool refract)
+ccl_device void svm_node_glass_setup(ccl_private ShaderData *sd,
+                                     ccl_private MicrofacetBsdf *bsdf,
+                                     int type,
+                                     float eta,
+                                     float roughness,
+                                     bool refract)
 {
   if (type == CLOSURE_BSDF_SHARP_GLASS_ID) {
     if (refract) {
@@ -58,8 +62,12 @@ ccl_device void svm_node_glass_setup(
 }
 
 template<uint node_feature_mask, ShaderType shader_type>
-ccl_device_noinline int svm_node_closure_bsdf(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int path_flag, int offset)
+ccl_device_noinline int svm_node_closure_bsdf(ccl_global const KernelGlobals *kg,
+                                              ccl_private ShaderData *sd,
+                                              ccl_private float *stack,
+                                              uint4 node,
+                                              int path_flag,
+                                              int offset)
 {
   uint type, param1_offset, param2_offset;
 
@@ -213,8 +221,8 @@ ccl_device_noinline int svm_node_closure_bsdf(
         if (subsurface <= CLOSURE_WEIGHT_CUTOFF && diffuse_weight > CLOSURE_WEIGHT_CUTOFF) {
           float3 diff_weight = weight * base_color * diffuse_weight;
 
-          PrincipledDiffuseBsdf *bsdf = (PrincipledDiffuseBsdf *)bsdf_alloc(
-              sd, sizeof(PrincipledDiffuseBsdf), diff_weight);
+          ccl_private PrincipledDiffuseBsdf *bsdf = (ccl_private PrincipledDiffuseBsdf *)
+              bsdf_alloc(sd, sizeof(PrincipledDiffuseBsdf), diff_weight);
 
           if (bsdf) {
             bsdf->N = N;
@@ -225,7 +233,7 @@ ccl_device_noinline int svm_node_closure_bsdf(
           }
         }
         else if (subsurface > CLOSURE_WEIGHT_CUTOFF) {
-          Bssrdf *bssrdf = bssrdf_alloc(sd, subsurf_weight);
+          ccl_private Bssrdf *bssrdf = bssrdf_alloc(sd, subsurf_weight);
 
           if (bssrdf) {
             bssrdf->radius = subsurface_radius * subsurface;
@@ -247,7 +255,7 @@ ccl_device_noinline int svm_node_closure_bsdf(
       if (diffuse_weight > CLOSURE_WEIGHT_CUTOFF) {
         float3 diff_weight = weight * base_color * diffuse_weight;
 
-        PrincipledDiffuseBsdf *bsdf = (PrincipledDiffuseBsdf *)bsdf_alloc(
+        ccl_private PrincipledDiffuseBsdf *bsdf = (ccl_private PrincipledDiffuseBsdf *)bsdf_alloc(
             sd, sizeof(PrincipledDiffuseBsdf), diff_weight);
 
         if (bsdf) {
@@ -273,7 +281,7 @@ ccl_device_noinline int svm_node_closure_bsdf(
 
         float3 sheen_weight = weight * sheen * sheen_color * diffuse_weight;
 
-        PrincipledSheenBsdf *bsdf = (PrincipledSheenBsdf *)bsdf_alloc(
+        ccl_private PrincipledSheenBsdf *bsdf = (ccl_private PrincipledSheenBsdf *)bsdf_alloc(
             sd, sizeof(PrincipledSheenBsdf), sheen_weight);
 
         if (bsdf) {
@@ -292,11 +300,12 @@ ccl_device_noinline int svm_node_closure_bsdf(
             (specular > CLOSURE_WEIGHT_CUTOFF || metallic > CLOSURE_WEIGHT_CUTOFF)) {
           float3 spec_weight = weight * specular_weight;
 
-          MicrofacetBsdf *bsdf = (MicrofacetBsdf *)bsdf_alloc(
+          ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
               sd, sizeof(MicrofacetBsdf), spec_weight);
-          MicrofacetExtra *extra = (bsdf != NULL) ? (MicrofacetExtra *)closure_alloc_extra(
-                                                        sd, sizeof(MicrofacetExtra)) :
-                                                    NULL;
+          ccl_private MicrofacetExtra *extra =
+              (bsdf != NULL) ?
+                  (ccl_private MicrofacetExtra *)closure_alloc_extra(sd, sizeof(MicrofacetExtra)) :
+                  NULL;
 
           if (bsdf && extra) {
             bsdf->N = N;
@@ -355,11 +364,12 @@ ccl_device_noinline int svm_node_closure_bsdf(
             if (kernel_data.integrator.caustics_reflective || (path_flag & PATH_RAY_DIFFUSE) == 0)
 #  endif
             {
-              MicrofacetBsdf *bsdf = (MicrofacetBsdf *)bsdf_alloc(
+              ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
                   sd, sizeof(MicrofacetBsdf), glass_weight * fresnel);
-              MicrofacetExtra *extra = (bsdf != NULL) ? (MicrofacetExtra *)closure_alloc_extra(
-                                                            sd, sizeof(MicrofacetExtra)) :
-                                                        NULL;
+              ccl_private MicrofacetExtra *extra =
+                  (bsdf != NULL) ? (ccl_private MicrofacetExtra *)closure_alloc_extra(
+                                       sd, sizeof(MicrofacetExtra)) :
+                                   NULL;
 
               if (bsdf && extra) {
                 bsdf->N = N;
@@ -384,7 +394,7 @@ ccl_device_noinline int svm_node_closure_bsdf(
             if (kernel_data.integrator.caustics_refractive || (path_flag & PATH_RAY_DIFFUSE) == 0)
 #  endif
             {
-              MicrofacetBsdf *bsdf = (MicrofacetBsdf *)bsdf_alloc(
+              ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
                   sd, sizeof(MicrofacetBsdf), base_color * glass_weight * (1.0f - fresnel));
               if (bsdf) {
                 bsdf->N = N;
@@ -407,11 +417,12 @@ ccl_device_noinline int svm_node_closure_bsdf(
             }
           }
           else { /* use multi-scatter GGX */
-            MicrofacetBsdf *bsdf = (MicrofacetBsdf *)bsdf_alloc(
+            ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
                 sd, sizeof(MicrofacetBsdf), glass_weight);
-            MicrofacetExtra *extra = (bsdf != NULL) ? (MicrofacetExtra *)closure_alloc_extra(
-                                                          sd, sizeof(MicrofacetExtra)) :
-                                                      NULL;
+            ccl_private MicrofacetExtra *extra =
+                (bsdf != NULL) ? (ccl_private MicrofacetExtra *)closure_alloc_extra(
+                                     sd, sizeof(MicrofacetExtra)) :
+                                 NULL;
 
             if (bsdf && extra) {
               bsdf->N = N;
@@ -440,10 +451,12 @@ ccl_device_noinline int svm_node_closure_bsdf(
       if (kernel_data.integrator.caustics_reflective || (path_flag & PATH_RAY_DIFFUSE) == 0) {
 #  endif
         if (clearcoat > CLOSURE_WEIGHT_CUTOFF) {
-          MicrofacetBsdf *bsdf = (MicrofacetBsdf *)bsdf_alloc(sd, sizeof(MicrofacetBsdf), weight);
-          MicrofacetExtra *extra = (bsdf != NULL) ? (MicrofacetExtra *)closure_alloc_extra(
-                                                        sd, sizeof(MicrofacetExtra)) :
-                                                    NULL;
+          ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
+              sd, sizeof(MicrofacetBsdf), weight);
+          ccl_private MicrofacetExtra *extra =
+              (bsdf != NULL) ?
+                  (ccl_private MicrofacetExtra *)closure_alloc_extra(sd, sizeof(MicrofacetExtra)) :
+                  NULL;
 
           if (bsdf && extra) {
             bsdf->N = clearcoat_normal;
@@ -471,7 +484,8 @@ ccl_device_noinline int svm_node_closure_bsdf(
 #endif /* __PRINCIPLED__ */
     case CLOSURE_BSDF_DIFFUSE_ID: {
       float3 weight = sd->svm_closure_weight * mix_weight;
-      OrenNayarBsdf *bsdf = (OrenNayarBsdf *)bsdf_alloc(sd, sizeof(OrenNayarBsdf), weight);
+      ccl_private OrenNayarBsdf *bsdf = (ccl_private OrenNayarBsdf *)bsdf_alloc(
+          sd, sizeof(OrenNayarBsdf), weight);
 
       if (bsdf) {
         bsdf->N = N;
@@ -479,7 +493,7 @@ ccl_device_noinline int svm_node_closure_bsdf(
         float roughness = param1;
 
         if (roughness == 0.0f) {
-          sd->flag |= bsdf_diffuse_setup((DiffuseBsdf *)bsdf);
+          sd->flag |= bsdf_diffuse_setup((ccl_private DiffuseBsdf *)bsdf);
         }
         else {
           bsdf->roughness = roughness;
@@ -490,7 +504,8 @@ ccl_device_noinline int svm_node_closure_bsdf(
     }
     case CLOSURE_BSDF_TRANSLUCENT_ID: {
       float3 weight = sd->svm_closure_weight * mix_weight;
-      DiffuseBsdf *bsdf = (DiffuseBsdf *)bsdf_alloc(sd, sizeof(DiffuseBsdf), weight);
+      ccl_private DiffuseBsdf *bsdf = (ccl_private DiffuseBsdf *)bsdf_alloc(
+          sd, sizeof(DiffuseBsdf), weight);
 
       if (bsdf) {
         bsdf->N = N;
@@ -513,7 +528,8 @@ ccl_device_noinline int svm_node_closure_bsdf(
         break;
 #endif
       float3 weight = sd->svm_closure_weight * mix_weight;
-      MicrofacetBsdf *bsdf = (MicrofacetBsdf *)bsdf_alloc(sd, sizeof(MicrofacetBsdf), weight);
+      ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
+          sd, sizeof(MicrofacetBsdf), weight);
 
       if (!bsdf) {
         break;
@@ -559,7 +575,8 @@ ccl_device_noinline int svm_node_closure_bsdf(
         sd->flag |= bsdf_microfacet_ggx_setup(bsdf);
       else if (type == CLOSURE_BSDF_MICROFACET_MULTI_GGX_ID) {
         kernel_assert(stack_valid(data_node.w));
-        bsdf->extra = (MicrofacetExtra *)closure_alloc_extra(sd, sizeof(MicrofacetExtra));
+        bsdf->extra = (ccl_private MicrofacetExtra *)closure_alloc_extra(sd,
+                                                                         sizeof(MicrofacetExtra));
         if (bsdf->extra) {
           bsdf->extra->color = stack_load_float3(stack, data_node.w);
           bsdf->extra->cspec0 = make_float3(0.0f, 0.0f, 0.0f);
@@ -581,7 +598,8 @@ ccl_device_noinline int svm_node_closure_bsdf(
         break;
 #endif
       float3 weight = sd->svm_closure_weight * mix_weight;
-      MicrofacetBsdf *bsdf = (MicrofacetBsdf *)bsdf_alloc(sd, sizeof(MicrofacetBsdf), weight);
+      ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
+          sd, sizeof(MicrofacetBsdf), weight);
 
       if (bsdf) {
         bsdf->N = N;
@@ -639,7 +657,7 @@ ccl_device_noinline int svm_node_closure_bsdf(
       if (kernel_data.integrator.caustics_reflective || (path_flag & PATH_RAY_DIFFUSE) == 0)
 #endif
       {
-        MicrofacetBsdf *bsdf = (MicrofacetBsdf *)bsdf_alloc(
+        ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
             sd, sizeof(MicrofacetBsdf), weight * fresnel);
 
         if (bsdf) {
@@ -655,7 +673,7 @@ ccl_device_noinline int svm_node_closure_bsdf(
       if (kernel_data.integrator.caustics_refractive || (path_flag & PATH_RAY_DIFFUSE) == 0)
 #endif
       {
-        MicrofacetBsdf *bsdf = (MicrofacetBsdf *)bsdf_alloc(
+        ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
             sd, sizeof(MicrofacetBsdf), weight * (1.0f - fresnel));
 
         if (bsdf) {
@@ -675,12 +693,14 @@ ccl_device_noinline int svm_node_closure_bsdf(
         break;
 #endif
       float3 weight = sd->svm_closure_weight * mix_weight;
-      MicrofacetBsdf *bsdf = (MicrofacetBsdf *)bsdf_alloc(sd, sizeof(MicrofacetBsdf), weight);
+      ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
+          sd, sizeof(MicrofacetBsdf), weight);
       if (!bsdf) {
         break;
       }
 
-      MicrofacetExtra *extra = (MicrofacetExtra *)closure_alloc_extra(sd, sizeof(MicrofacetExtra));
+      ccl_private MicrofacetExtra *extra = (ccl_private MicrofacetExtra *)closure_alloc_extra(
+          sd, sizeof(MicrofacetExtra));
       if (!extra) {
         break;
       }
@@ -706,7 +726,8 @@ ccl_device_noinline int svm_node_closure_bsdf(
     }
     case CLOSURE_BSDF_ASHIKHMIN_VELVET_ID: {
       float3 weight = sd->svm_closure_weight * mix_weight;
-      VelvetBsdf *bsdf = (VelvetBsdf *)bsdf_alloc(sd, sizeof(VelvetBsdf), weight);
+      ccl_private VelvetBsdf *bsdf = (ccl_private VelvetBsdf *)bsdf_alloc(
+          sd, sizeof(VelvetBsdf), weight);
 
       if (bsdf) {
         bsdf->N = N;
@@ -724,7 +745,8 @@ ccl_device_noinline int svm_node_closure_bsdf(
 #endif
     case CLOSURE_BSDF_DIFFUSE_TOON_ID: {
       float3 weight = sd->svm_closure_weight * mix_weight;
-      ToonBsdf *bsdf = (ToonBsdf *)bsdf_alloc(sd, sizeof(ToonBsdf), weight);
+      ccl_private ToonBsdf *bsdf = (ccl_private ToonBsdf *)bsdf_alloc(
+          sd, sizeof(ToonBsdf), weight);
 
       if (bsdf) {
         bsdf->N = N;
@@ -771,11 +793,11 @@ ccl_device_noinline int svm_node_closure_bsdf(
         random = stack_load_float_default(stack, random_ofs, data_node3.y);
       }
 
-      PrincipledHairBSDF *bsdf = (PrincipledHairBSDF *)bsdf_alloc(
+      ccl_private PrincipledHairBSDF *bsdf = (ccl_private PrincipledHairBSDF *)bsdf_alloc(
           sd, sizeof(PrincipledHairBSDF), weight);
       if (bsdf) {
-        PrincipledHairExtra *extra = (PrincipledHairExtra *)closure_alloc_extra(
-            sd, sizeof(PrincipledHairExtra));
+        ccl_private PrincipledHairExtra *extra = (ccl_private PrincipledHairExtra *)
+            closure_alloc_extra(sd, sizeof(PrincipledHairExtra));
 
         if (!extra)
           break;
@@ -854,7 +876,8 @@ ccl_device_noinline int svm_node_closure_bsdf(
     case CLOSURE_BSDF_HAIR_TRANSMISSION_ID: {
       float3 weight = sd->svm_closure_weight * mix_weight;
 
-      HairBsdf *bsdf = (HairBsdf *)bsdf_alloc(sd, sizeof(HairBsdf), weight);
+      ccl_private HairBsdf *bsdf = (ccl_private HairBsdf *)bsdf_alloc(
+          sd, sizeof(HairBsdf), weight);
 
       if (bsdf) {
         bsdf->N = N;
@@ -889,7 +912,7 @@ ccl_device_noinline int svm_node_closure_bsdf(
     case CLOSURE_BSSRDF_RANDOM_WALK_ID:
     case CLOSURE_BSSRDF_RANDOM_WALK_FIXED_RADIUS_ID: {
       float3 weight = sd->svm_closure_weight * mix_weight;
-      Bssrdf *bssrdf = bssrdf_alloc(sd, weight);
+      ccl_private Bssrdf *bssrdf = bssrdf_alloc(sd, weight);
 
       if (bssrdf) {
         /* disable in case of diffuse ancestor, can't see it well then and
@@ -921,9 +944,9 @@ ccl_device_noinline int svm_node_closure_bsdf(
 }
 
 template<ShaderType shader_type>
-ccl_device_noinline void svm_node_closure_volume(const KernelGlobals *kg,
-                                                 ShaderData *sd,
-                                                 float *stack,
+ccl_device_noinline void svm_node_closure_volume(ccl_global const KernelGlobals *kg,
+                                                 ccl_private ShaderData *sd,
+                                                 ccl_private float *stack,
                                                  uint4 node)
 {
 #ifdef __VOLUME__
@@ -958,7 +981,7 @@ ccl_device_noinline void svm_node_closure_volume(const KernelGlobals *kg,
 
   /* Add closure for volume scattering. */
   if (type == CLOSURE_VOLUME_HENYEY_GREENSTEIN_ID) {
-    HenyeyGreensteinVolume *volume = (HenyeyGreensteinVolume *)bsdf_alloc(
+    ccl_private HenyeyGreensteinVolume *volume = (ccl_private HenyeyGreensteinVolume *)bsdf_alloc(
         sd, sizeof(HenyeyGreensteinVolume), weight);
 
     if (volume) {
@@ -976,8 +999,12 @@ ccl_device_noinline void svm_node_closure_volume(const KernelGlobals *kg,
 }
 
 template<ShaderType shader_type>
-ccl_device_noinline int svm_node_principled_volume(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int path_flag, int offset)
+ccl_device_noinline int svm_node_principled_volume(ccl_global const KernelGlobals *kg,
+                                                   ccl_private ShaderData *sd,
+                                                   ccl_private float *stack,
+                                                   uint4 node,
+                                                   int path_flag,
+                                                   int offset)
 {
 #ifdef __VOLUME__
   uint4 value_node = read_node(kg, &offset);
@@ -1023,7 +1050,7 @@ ccl_device_noinline int svm_node_principled_volume(
     }
 
     /* Add closure for volume scattering. */
-    HenyeyGreensteinVolume *volume = (HenyeyGreensteinVolume *)bsdf_alloc(
+    ccl_private HenyeyGreensteinVolume *volume = (ccl_private HenyeyGreensteinVolume *)bsdf_alloc(
         sd, sizeof(HenyeyGreensteinVolume), color * density);
     if (volume) {
       float anisotropy = (stack_valid(anisotropy_offset)) ?
@@ -1087,7 +1114,9 @@ ccl_device_noinline int svm_node_principled_volume(
   return offset;
 }
 
-ccl_device_noinline void svm_node_closure_emission(ShaderData *sd, float *stack, uint4 node)
+ccl_device_noinline void svm_node_closure_emission(ccl_private ShaderData *sd,
+                                                   ccl_private float *stack,
+                                                   uint4 node)
 {
   uint mix_weight_offset = node.y;
   float3 weight = sd->svm_closure_weight;
@@ -1104,7 +1133,9 @@ ccl_device_noinline void svm_node_closure_emission(ShaderData *sd, float *stack,
   emission_setup(sd, weight);
 }
 
-ccl_device_noinline void svm_node_closure_background(ShaderData *sd, float *stack, uint4 node)
+ccl_device_noinline void svm_node_closure_background(ccl_private ShaderData *sd,
+                                                     ccl_private float *stack,
+                                                     uint4 node)
 {
   uint mix_weight_offset = node.y;
   float3 weight = sd->svm_closure_weight;
@@ -1121,7 +1152,9 @@ ccl_device_noinline void svm_node_closure_background(ShaderData *sd, float *stac
   background_setup(sd, weight);
 }
 
-ccl_device_noinline void svm_node_closure_holdout(ShaderData *sd, float *stack, uint4 node)
+ccl_device_noinline void svm_node_closure_holdout(ccl_private ShaderData *sd,
+                                                  ccl_private float *stack,
+                                                  uint4 node)
 {
   uint mix_weight_offset = node.y;
 
@@ -1142,26 +1175,28 @@ ccl_device_noinline void svm_node_closure_holdout(ShaderData *sd, float *stack,
 
 /* Closure Nodes */
 
-ccl_device_inline void svm_node_closure_store_weight(ShaderData *sd, float3 weight)
+ccl_device_inline void svm_node_closure_store_weight(ccl_private ShaderData *sd, float3 weight)
 {
   sd->svm_closure_weight = weight;
 }
 
-ccl_device void svm_node_closure_set_weight(ShaderData *sd, uint r, uint g, uint b)
+ccl_device void svm_node_closure_set_weight(ccl_private ShaderData *sd, uint r, uint g, uint b)
 {
   float3 weight = make_float3(__uint_as_float(r), __uint_as_float(g), __uint_as_float(b));
   svm_node_closure_store_weight(sd, weight);
 }
 
-ccl_device void svm_node_closure_weight(ShaderData *sd, float *stack, uint weight_offset)
+ccl_device void svm_node_closure_weight(ccl_private ShaderData *sd,
+                                        ccl_private float *stack,
+                                        uint weight_offset)
 {
   float3 weight = stack_load_float3(stack, weight_offset);
   svm_node_closure_store_weight(sd, weight);
 }
 
-ccl_device_noinline void svm_node_emission_weight(const KernelGlobals *kg,
-                                                  ShaderData *sd,
-                                                  float *stack,
+ccl_device_noinline void svm_node_emission_weight(ccl_global const KernelGlobals *kg,
+                                                  ccl_private ShaderData *sd,
+                                                  ccl_private float *stack,
                                                   uint4 node)
 {
   uint color_offset = node.y;
@@ -1173,7 +1208,9 @@ ccl_device_noinline void svm_node_emission_weight(const KernelGlobals *kg,
   svm_node_closure_store_weight(sd, weight);
 }
 
-ccl_device_noinline void svm_node_mix_closure(ShaderData *sd, float *stack, uint4 node)
+ccl_device_noinline void svm_node_mix_closure(ccl_private ShaderData *sd,
+                                              ccl_private float *stack,
+                                              uint4 node)
 {
   /* fetch weight from blend input, previous mix closures,
    * and write to stack to be used by closure nodes later */
@@ -1195,8 +1232,11 @@ ccl_device_noinline void svm_node_mix_closure(ShaderData *sd, float *stack, uint
 
 /* (Bump) normal */
 
-ccl_device void svm_node_set_normal(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint in_direction, uint out_normal)
+ccl_device void svm_node_set_normal(ccl_global const KernelGlobals *kg,
+                                    ccl_private ShaderData *sd,
+                                    ccl_private float *stack,
+                                    uint in_direction,
+                                    uint out_normal)
 {
   float3 normal = stack_load_float3(stack, in_direction);
   sd->N = normal;
diff --git a/intern/cycles/kernel/svm/svm_convert.h b/intern/cycles/kernel/svm/svm_convert.h
index 37d40167ccc..0d53779a5c8 100644
--- a/intern/cycles/kernel/svm/svm_convert.h
+++ b/intern/cycles/kernel/svm/svm_convert.h
@@ -18,8 +18,12 @@ CCL_NAMESPACE_BEGIN
 
 /* Conversion Nodes */
 
-ccl_device_noinline void svm_node_convert(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint from, uint to)
+ccl_device_noinline void svm_node_convert(ccl_global const KernelGlobals *kg,
+                                          ccl_private ShaderData *sd,
+                                          ccl_private float *stack,
+                                          uint type,
+                                          uint from,
+                                          uint to)
 {
   switch (type) {
     case NODE_CONVERT_FI: {
diff --git a/intern/cycles/kernel/svm/svm_displace.h b/intern/cycles/kernel/svm/svm_displace.h
index a1d952173d8..7a3c8a6d36d 100644
--- a/intern/cycles/kernel/svm/svm_displace.h
+++ b/intern/cycles/kernel/svm/svm_displace.h
@@ -20,9 +20,9 @@ CCL_NAMESPACE_BEGIN
 
 /* Bump Node */
 
-ccl_device_noinline void svm_node_set_bump(const KernelGlobals *kg,
-                                           ShaderData *sd,
-                                           float *stack,
+ccl_device_noinline void svm_node_set_bump(ccl_global const KernelGlobals *kg,
+                                           ccl_private ShaderData *sd,
+                                           ccl_private float *stack,
                                            uint4 node)
 {
 #ifdef __RAY_DIFFERENTIALS__
@@ -88,18 +88,18 @@ ccl_device_noinline void svm_node_set_bump(const KernelGlobals *kg,
 
 /* Displacement Node */
 
-ccl_device void svm_node_set_displacement(const KernelGlobals *kg,
-                                          ShaderData *sd,
-                                          float *stack,
+ccl_device void svm_node_set_displacement(ccl_global const KernelGlobals *kg,
+                                          ccl_private ShaderData *sd,
+                                          ccl_private float *stack,
                                           uint fac_offset)
 {
   float3 dP = stack_load_float3(stack, fac_offset);
   sd->P += dP;
 }
 
-ccl_device_noinline void svm_node_displacement(const KernelGlobals *kg,
-                                               ShaderData *sd,
-                                               float *stack,
+ccl_device_noinline void svm_node_displacement(ccl_global const KernelGlobals *kg,
+                                               ccl_private ShaderData *sd,
+                                               ccl_private float *stack,
                                                uint4 node)
 {
   uint height_offset, midlevel_offset, scale_offset, normal_offset;
@@ -127,8 +127,11 @@ ccl_device_noinline void svm_node_displacement(const KernelGlobals *kg,
   stack_store_float3(stack, node.z, dP);
 }
 
-ccl_device_noinline int svm_node_vector_displacement(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_vector_displacement(ccl_global const KernelGlobals *kg,
+                                                     ccl_private ShaderData *sd,
+                                                     ccl_private float *stack,
+                                                     uint4 node,
+                                                     int offset)
 {
   uint4 data_node = read_node(kg, &offset);
   uint space = data_node.x;
diff --git a/intern/cycles/kernel/svm/svm_fresnel.h b/intern/cycles/kernel/svm/svm_fresnel.h
index b5ecdbe2abf..449ec84370f 100644
--- a/intern/cycles/kernel/svm/svm_fresnel.h
+++ b/intern/cycles/kernel/svm/svm_fresnel.h
@@ -18,8 +18,11 @@ CCL_NAMESPACE_BEGIN
 
 /* Fresnel Node */
 
-ccl_device_noinline void svm_node_fresnel(
-    ShaderData *sd, float *stack, uint ior_offset, uint ior_value, uint node)
+ccl_device_noinline void svm_node_fresnel(ccl_private ShaderData *sd,
+                                          ccl_private float *stack,
+                                          uint ior_offset,
+                                          uint ior_value,
+                                          uint node)
 {
   uint normal_offset, out_offset;
   svm_unpack_node_uchar2(node, &normal_offset, &out_offset);
@@ -37,7 +40,9 @@ ccl_device_noinline void svm_node_fresnel(
 
 /* Layer Weight Node */
 
-ccl_device_noinline void svm_node_layer_weight(ShaderData *sd, float *stack, uint4 node)
+ccl_device_noinline void svm_node_layer_weight(ccl_private ShaderData *sd,
+                                               ccl_private float *stack,
+                                               uint4 node)
 {
   uint blend_offset = node.y;
   uint blend_value = node.z;
diff --git a/intern/cycles/kernel/svm/svm_gamma.h b/intern/cycles/kernel/svm/svm_gamma.h
index f6fafdee941..7ec6c31065d 100644
--- a/intern/cycles/kernel/svm/svm_gamma.h
+++ b/intern/cycles/kernel/svm/svm_gamma.h
@@ -16,8 +16,11 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device_noinline void svm_node_gamma(
-    ShaderData *sd, float *stack, uint in_gamma, uint in_color, uint out_color)
+ccl_device_noinline void svm_node_gamma(ccl_private ShaderData *sd,
+                                        ccl_private float *stack,
+                                        uint in_gamma,
+                                        uint in_color,
+                                        uint out_color)
 {
   float3 color = stack_load_float3(stack, in_color);
   float gamma = stack_load_float(stack, in_gamma);
diff --git a/intern/cycles/kernel/svm/svm_geometry.h b/intern/cycles/kernel/svm/svm_geometry.h
index 432529eb061..a94464d3a52 100644
--- a/intern/cycles/kernel/svm/svm_geometry.h
+++ b/intern/cycles/kernel/svm/svm_geometry.h
@@ -18,8 +18,11 @@ CCL_NAMESPACE_BEGIN
 
 /* Geometry Node */
 
-ccl_device_noinline void svm_node_geometry(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset)
+ccl_device_noinline void svm_node_geometry(ccl_global const KernelGlobals *kg,
+                                           ccl_private ShaderData *sd,
+                                           ccl_private float *stack,
+                                           uint type,
+                                           uint out_offset)
 {
   float3 data;
 
@@ -51,8 +54,11 @@ ccl_device_noinline void svm_node_geometry(
   stack_store_float3(stack, out_offset, data);
 }
 
-ccl_device_noinline void svm_node_geometry_bump_dx(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset)
+ccl_device_noinline void svm_node_geometry_bump_dx(ccl_global const KernelGlobals *kg,
+                                                   ccl_private ShaderData *sd,
+                                                   ccl_private float *stack,
+                                                   uint type,
+                                                   uint out_offset)
 {
 #ifdef __RAY_DIFFERENTIALS__
   float3 data;
@@ -75,8 +81,11 @@ ccl_device_noinline void svm_node_geometry_bump_dx(
 #endif
 }
 
-ccl_device_noinline void svm_node_geometry_bump_dy(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset)
+ccl_device_noinline void svm_node_geometry_bump_dy(ccl_global const KernelGlobals *kg,
+                                                   ccl_private ShaderData *sd,
+                                                   ccl_private float *stack,
+                                                   uint type,
+                                                   uint out_offset)
 {
 #ifdef __RAY_DIFFERENTIALS__
   float3 data;
@@ -101,8 +110,11 @@ ccl_device_noinline void svm_node_geometry_bump_dy(
 
 /* Object Info */
 
-ccl_device_noinline void svm_node_object_info(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset)
+ccl_device_noinline void svm_node_object_info(ccl_global const KernelGlobals *kg,
+                                              ccl_private ShaderData *sd,
+                                              ccl_private float *stack,
+                                              uint type,
+                                              uint out_offset)
 {
   float data;
 
@@ -140,8 +152,11 @@ ccl_device_noinline void svm_node_object_info(
 
 /* Particle Info */
 
-ccl_device_noinline void svm_node_particle_info(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset)
+ccl_device_noinline void svm_node_particle_info(ccl_global const KernelGlobals *kg,
+                                                ccl_private ShaderData *sd,
+                                                ccl_private float *stack,
+                                                uint type,
+                                                uint out_offset)
 {
   switch (type) {
     case NODE_INFO_PAR_INDEX: {
@@ -199,8 +214,11 @@ ccl_device_noinline void svm_node_particle_info(
 
 /* Hair Info */
 
-ccl_device_noinline void svm_node_hair_info(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset)
+ccl_device_noinline void svm_node_hair_info(ccl_global const KernelGlobals *kg,
+                                            ccl_private ShaderData *sd,
+                                            ccl_private float *stack,
+                                            uint type,
+                                            uint out_offset)
 {
   float data;
   float3 data3;
diff --git a/intern/cycles/kernel/svm/svm_gradient.h b/intern/cycles/kernel/svm/svm_gradient.h
index cd15f7097e7..8cc37be606f 100644
--- a/intern/cycles/kernel/svm/svm_gradient.h
+++ b/intern/cycles/kernel/svm/svm_gradient.h
@@ -60,7 +60,9 @@ ccl_device float svm_gradient(float3 p, NodeGradientType type)
   return 0.0f;
 }
 
-ccl_device_noinline void svm_node_tex_gradient(ShaderData *sd, float *stack, uint4 node)
+ccl_device_noinline void svm_node_tex_gradient(ccl_private ShaderData *sd,
+                                               ccl_private float *stack,
+                                               uint4 node)
 {
   uint type, co_offset, color_offset, fac_offset;
 
diff --git a/intern/cycles/kernel/svm/svm_hsv.h b/intern/cycles/kernel/svm/svm_hsv.h
index 6f49a8385aa..feb85eda122 100644
--- a/intern/cycles/kernel/svm/svm_hsv.h
+++ b/intern/cycles/kernel/svm/svm_hsv.h
@@ -19,9 +19,9 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device_noinline void svm_node_hsv(const KernelGlobals *kg,
-                                      ShaderData *sd,
-                                      float *stack,
+ccl_device_noinline void svm_node_hsv(ccl_global const KernelGlobals *kg,
+                                      ccl_private ShaderData *sd,
+                                      ccl_private float *stack,
                                       uint4 node)
 {
   uint in_color_offset, fac_offset, out_color_offset;
diff --git a/intern/cycles/kernel/svm/svm_ies.h b/intern/cycles/kernel/svm/svm_ies.h
index 9c13734ecf0..7d41205c9ef 100644
--- a/intern/cycles/kernel/svm/svm_ies.h
+++ b/intern/cycles/kernel/svm/svm_ies.h
@@ -19,7 +19,7 @@ CCL_NAMESPACE_BEGIN
 /* IES Light */
 
 ccl_device_inline float interpolate_ies_vertical(
-    const KernelGlobals *kg, int ofs, int v, int v_num, float v_frac, int h)
+    ccl_global const KernelGlobals *kg, int ofs, int v, int v_num, float v_frac, int h)
 {
   /* Since lookups are performed in spherical coordinates, clamping the coordinates at the low end
    * of v (corresponding to the north pole) would result in artifacts. The proper way of dealing
@@ -39,7 +39,7 @@ ccl_device_inline float interpolate_ies_vertical(
   return cubic_interp(a, b, c, d, v_frac);
 }
 
-ccl_device_inline float kernel_ies_interp(const KernelGlobals *kg,
+ccl_device_inline float kernel_ies_interp(ccl_global const KernelGlobals *kg,
                                           int slot,
                                           float h_angle,
                                           float v_angle)
@@ -98,9 +98,9 @@ ccl_device_inline float kernel_ies_interp(const KernelGlobals *kg,
   return max(cubic_interp(a, b, c, d, h_frac), 0.0f);
 }
 
-ccl_device_noinline void svm_node_ies(const KernelGlobals *kg,
-                                      ShaderData *sd,
-                                      float *stack,
+ccl_device_noinline void svm_node_ies(ccl_global const KernelGlobals *kg,
+                                      ccl_private ShaderData *sd,
+                                      ccl_private float *stack,
                                       uint4 node)
 {
   uint vector_offset, strength_offset, fac_offset, slot = node.z;
diff --git a/intern/cycles/kernel/svm/svm_image.h b/intern/cycles/kernel/svm/svm_image.h
index ce70109392b..2de80d5fc29 100644
--- a/intern/cycles/kernel/svm/svm_image.h
+++ b/intern/cycles/kernel/svm/svm_image.h
@@ -16,7 +16,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device float4 svm_image_texture(const KernelGlobals *kg, int id, float x, float y, uint flags)
+ccl_device float4
+svm_image_texture(ccl_global const KernelGlobals *kg, int id, float x, float y, uint flags)
 {
   if (id == -1) {
     return make_float4(
@@ -44,8 +45,11 @@ ccl_device_inline float3 texco_remap_square(float3 co)
   return (co - make_float3(0.5f, 0.5f, 0.5f)) * 2.0f;
 }
 
-ccl_device_noinline int svm_node_tex_image(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_tex_image(ccl_global const KernelGlobals *kg,
+                                           ccl_private ShaderData *sd,
+                                           ccl_private float *stack,
+                                           uint4 node,
+                                           int offset)
 {
   uint co_offset, out_offset, alpha_offset, flags;
 
@@ -117,9 +121,9 @@ ccl_device_noinline int svm_node_tex_image(
   return offset;
 }
 
-ccl_device_noinline void svm_node_tex_image_box(const KernelGlobals *kg,
-                                                ShaderData *sd,
-                                                float *stack,
+ccl_device_noinline void svm_node_tex_image_box(ccl_global const KernelGlobals *kg,
+                                                ccl_private ShaderData *sd,
+                                                ccl_private float *stack,
                                                 uint4 node)
 {
   /* get object space normal */
@@ -219,9 +223,9 @@ ccl_device_noinline void svm_node_tex_image_box(const KernelGlobals *kg,
     stack_store_float(stack, alpha_offset, f.w);
 }
 
-ccl_device_noinline void svm_node_tex_environment(const KernelGlobals *kg,
-                                                  ShaderData *sd,
-                                                  float *stack,
+ccl_device_noinline void svm_node_tex_environment(ccl_global const KernelGlobals *kg,
+                                                  ccl_private ShaderData *sd,
+                                                  ccl_private float *stack,
                                                   uint4 node)
 {
   uint id = node.y;
diff --git a/intern/cycles/kernel/svm/svm_invert.h b/intern/cycles/kernel/svm/svm_invert.h
index 27cdaaff473..60668ec00f1 100644
--- a/intern/cycles/kernel/svm/svm_invert.h
+++ b/intern/cycles/kernel/svm/svm_invert.h
@@ -21,8 +21,11 @@ ccl_device float invert(float color, float factor)
   return factor * (1.0f - color) + (1.0f - factor) * color;
 }
 
-ccl_device_noinline void svm_node_invert(
-    ShaderData *sd, float *stack, uint in_fac, uint in_color, uint out_color)
+ccl_device_noinline void svm_node_invert(ccl_private ShaderData *sd,
+                                         ccl_private float *stack,
+                                         uint in_fac,
+                                         uint in_color,
+                                         uint out_color)
 {
   float factor = stack_load_float(stack, in_fac);
   float3 color = stack_load_float3(stack, in_color);
diff --git a/intern/cycles/kernel/svm/svm_light_path.h b/intern/cycles/kernel/svm/svm_light_path.h
index 49fabad1cc5..aaff8376c7c 100644
--- a/intern/cycles/kernel/svm/svm_light_path.h
+++ b/intern/cycles/kernel/svm/svm_light_path.h
@@ -19,8 +19,8 @@ CCL_NAMESPACE_BEGIN
 /* Light Path Node */
 
 ccl_device_noinline void svm_node_light_path(INTEGRATOR_STATE_CONST_ARGS,
-                                             const ShaderData *sd,
-                                             float *stack,
+                                             ccl_private const ShaderData *sd,
+                                             ccl_private float *stack,
                                              uint type,
                                              uint out_offset,
                                              int path_flag)
@@ -106,7 +106,9 @@ ccl_device_noinline void svm_node_light_path(INTEGRATOR_STATE_CONST_ARGS,
 
 /* Light Falloff Node */
 
-ccl_device_noinline void svm_node_light_falloff(ShaderData *sd, float *stack, uint4 node)
+ccl_device_noinline void svm_node_light_falloff(ccl_private ShaderData *sd,
+                                                ccl_private float *stack,
+                                                uint4 node)
 {
   uint strength_offset, out_offset, smooth_offset;
 
diff --git a/intern/cycles/kernel/svm/svm_magic.h b/intern/cycles/kernel/svm/svm_magic.h
index 8784c760860..4c4f3bcf523 100644
--- a/intern/cycles/kernel/svm/svm_magic.h
+++ b/intern/cycles/kernel/svm/svm_magic.h
@@ -87,8 +87,11 @@ ccl_device_noinline_cpu float3 svm_magic(float3 p, int n, float distortion)
   return make_float3(0.5f - x, 0.5f - y, 0.5f - z);
 }
 
-ccl_device_noinline int svm_node_tex_magic(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_tex_magic(ccl_global const KernelGlobals *kg,
+                                           ccl_private ShaderData *sd,
+                                           ccl_private float *stack,
+                                           uint4 node,
+                                           int offset)
 {
   uint depth;
   uint scale_offset, distortion_offset, co_offset, fac_offset, color_offset;
diff --git a/intern/cycles/kernel/svm/svm_map_range.h b/intern/cycles/kernel/svm/svm_map_range.h
index c8684981e31..f4f7d3ca76f 100644
--- a/intern/cycles/kernel/svm/svm_map_range.h
+++ b/intern/cycles/kernel/svm/svm_map_range.h
@@ -24,9 +24,9 @@ ccl_device_inline float smootherstep(float edge0, float edge1, float x)
   return x * x * x * (x * (x * 6.0f - 15.0f) + 10.0f);
 }
 
-ccl_device_noinline int svm_node_map_range(const KernelGlobals *kg,
-                                           ShaderData *sd,
-                                           float *stack,
+ccl_device_noinline int svm_node_map_range(ccl_global const KernelGlobals *kg,
+                                           ccl_private ShaderData *sd,
+                                           ccl_private float *stack,
                                            uint value_stack_offset,
                                            uint parameters_stack_offsets,
                                            uint results_stack_offsets,
diff --git a/intern/cycles/kernel/svm/svm_mapping.h b/intern/cycles/kernel/svm/svm_mapping.h
index fcc724405f5..8102afc637e 100644
--- a/intern/cycles/kernel/svm/svm_mapping.h
+++ b/intern/cycles/kernel/svm/svm_mapping.h
@@ -18,9 +18,9 @@ CCL_NAMESPACE_BEGIN
 
 /* Mapping Node */
 
-ccl_device_noinline void svm_node_mapping(const KernelGlobals *kg,
-                                          ShaderData *sd,
-                                          float *stack,
+ccl_device_noinline void svm_node_mapping(ccl_global const KernelGlobals *kg,
+                                          ccl_private ShaderData *sd,
+                                          ccl_private float *stack,
                                           uint type,
                                           uint inputs_stack_offsets,
                                           uint result_stack_offset)
@@ -43,9 +43,9 @@ ccl_device_noinline void svm_node_mapping(const KernelGlobals *kg,
 
 /* Texture Mapping */
 
-ccl_device_noinline int svm_node_texture_mapping(const KernelGlobals *kg,
-                                                 ShaderData *sd,
-                                                 float *stack,
+ccl_device_noinline int svm_node_texture_mapping(ccl_global const KernelGlobals *kg,
+                                                 ccl_private ShaderData *sd,
+                                                 ccl_private float *stack,
                                                  uint vec_offset,
                                                  uint out_offset,
                                                  int offset)
@@ -62,9 +62,9 @@ ccl_device_noinline int svm_node_texture_mapping(const KernelGlobals *kg,
   return offset;
 }
 
-ccl_device_noinline int svm_node_min_max(const KernelGlobals *kg,
-                                         ShaderData *sd,
-                                         float *stack,
+ccl_device_noinline int svm_node_min_max(ccl_global const KernelGlobals *kg,
+                                         ccl_private ShaderData *sd,
+                                         ccl_private float *stack,
                                          uint vec_offset,
                                          uint out_offset,
                                          int offset)
diff --git a/intern/cycles/kernel/svm/svm_math.h b/intern/cycles/kernel/svm/svm_math.h
index 99e7a8f2bda..3897a453873 100644
--- a/intern/cycles/kernel/svm/svm_math.h
+++ b/intern/cycles/kernel/svm/svm_math.h
@@ -16,9 +16,9 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device_noinline void svm_node_math(const KernelGlobals *kg,
-                                       ShaderData *sd,
-                                       float *stack,
+ccl_device_noinline void svm_node_math(ccl_global const KernelGlobals *kg,
+                                       ccl_private ShaderData *sd,
+                                       ccl_private float *stack,
                                        uint type,
                                        uint inputs_stack_offsets,
                                        uint result_stack_offset)
@@ -34,9 +34,9 @@ ccl_device_noinline void svm_node_math(const KernelGlobals *kg,
   stack_store_float(stack, result_stack_offset, result);
 }
 
-ccl_device_noinline int svm_node_vector_math(const KernelGlobals *kg,
-                                             ShaderData *sd,
-                                             float *stack,
+ccl_device_noinline int svm_node_vector_math(ccl_global const KernelGlobals *kg,
+                                             ccl_private ShaderData *sd,
+                                             ccl_private float *stack,
                                              uint type,
                                              uint inputs_stack_offsets,
                                              uint outputs_stack_offsets,
diff --git a/intern/cycles/kernel/svm/svm_math_util.h b/intern/cycles/kernel/svm/svm_math_util.h
index 11b1e8f57f8..d3225b55ef0 100644
--- a/intern/cycles/kernel/svm/svm_math_util.h
+++ b/intern/cycles/kernel/svm/svm_math_util.h
@@ -16,8 +16,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device void svm_vector_math(float *value,
-                                float3 *vector,
+ccl_device void svm_vector_math(ccl_private float *value,
+                                ccl_private float3 *vector,
                                 NodeVectorMathType type,
                                 float3 a,
                                 float3 b,
diff --git a/intern/cycles/kernel/svm/svm_mix.h b/intern/cycles/kernel/svm/svm_mix.h
index 3e38080977f..0064c5e643c 100644
--- a/intern/cycles/kernel/svm/svm_mix.h
+++ b/intern/cycles/kernel/svm/svm_mix.h
@@ -18,9 +18,9 @@ CCL_NAMESPACE_BEGIN
 
 /* Node */
 
-ccl_device_noinline int svm_node_mix(const KernelGlobals *kg,
-                                     ShaderData *sd,
-                                     float *stack,
+ccl_device_noinline int svm_node_mix(ccl_global const KernelGlobals *kg,
+                                     ccl_private ShaderData *sd,
+                                     ccl_private float *stack,
                                      uint fac_offset,
                                      uint c1_offset,
                                      uint c2_offset,
diff --git a/intern/cycles/kernel/svm/svm_musgrave.h b/intern/cycles/kernel/svm/svm_musgrave.h
index 03a8b68b3ef..8523f45b95f 100644
--- a/intern/cycles/kernel/svm/svm_musgrave.h
+++ b/intern/cycles/kernel/svm/svm_musgrave.h
@@ -700,9 +700,9 @@ ccl_device_noinline_cpu float noise_musgrave_ridged_multi_fractal_4d(
   return value;
 }
 
-ccl_device_noinline int svm_node_tex_musgrave(const KernelGlobals *kg,
-                                              ShaderData *sd,
-                                              float *stack,
+ccl_device_noinline int svm_node_tex_musgrave(ccl_global const KernelGlobals *kg,
+                                              ccl_private ShaderData *sd,
+                                              ccl_private float *stack,
                                               uint offsets1,
                                               uint offsets2,
                                               uint offsets3,
diff --git a/intern/cycles/kernel/svm/svm_noisetex.h b/intern/cycles/kernel/svm/svm_noisetex.h
index 29b262ac06e..61da8227efa 100644
--- a/intern/cycles/kernel/svm/svm_noisetex.h
+++ b/intern/cycles/kernel/svm/svm_noisetex.h
@@ -55,8 +55,8 @@ ccl_device void noise_texture_1d(float co,
                                  float roughness,
                                  float distortion,
                                  bool color_is_needed,
-                                 float *value,
-                                 float3 *color)
+                                 ccl_private float *value,
+                                 ccl_private float3 *color)
 {
   float p = co;
   if (distortion != 0.0f) {
@@ -76,8 +76,8 @@ ccl_device void noise_texture_2d(float2 co,
                                  float roughness,
                                  float distortion,
                                  bool color_is_needed,
-                                 float *value,
-                                 float3 *color)
+                                 ccl_private float *value,
+                                 ccl_private float3 *color)
 {
   float2 p = co;
   if (distortion != 0.0f) {
@@ -98,8 +98,8 @@ ccl_device void noise_texture_3d(float3 co,
                                  float roughness,
                                  float distortion,
                                  bool color_is_needed,
-                                 float *value,
-                                 float3 *color)
+                                 ccl_private float *value,
+                                 ccl_private float3 *color)
 {
   float3 p = co;
   if (distortion != 0.0f) {
@@ -121,8 +121,8 @@ ccl_device void noise_texture_4d(float4 co,
                                  float roughness,
                                  float distortion,
                                  bool color_is_needed,
-                                 float *value,
-                                 float3 *color)
+                                 ccl_private float *value,
+                                 ccl_private float3 *color)
 {
   float4 p = co;
   if (distortion != 0.0f) {
@@ -140,9 +140,9 @@ ccl_device void noise_texture_4d(float4 co,
   }
 }
 
-ccl_device_noinline int svm_node_tex_noise(const KernelGlobals *kg,
-                                           ShaderData *sd,
-                                           float *stack,
+ccl_device_noinline int svm_node_tex_noise(ccl_global const KernelGlobals *kg,
+                                           ccl_private ShaderData *sd,
+                                           ccl_private float *stack,
                                            uint dimensions,
                                            uint offsets1,
                                            uint offsets2,
diff --git a/intern/cycles/kernel/svm/svm_normal.h b/intern/cycles/kernel/svm/svm_normal.h
index 724b5f281f9..0d1b4200d54 100644
--- a/intern/cycles/kernel/svm/svm_normal.h
+++ b/intern/cycles/kernel/svm/svm_normal.h
@@ -16,9 +16,9 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device_noinline int svm_node_normal(const KernelGlobals *kg,
-                                        ShaderData *sd,
-                                        float *stack,
+ccl_device_noinline int svm_node_normal(ccl_global const KernelGlobals *kg,
+                                        ccl_private ShaderData *sd,
+                                        ccl_private float *stack,
                                         uint in_normal_offset,
                                         uint out_normal_offset,
                                         uint out_dot_offset,
diff --git a/intern/cycles/kernel/svm/svm_ramp.h b/intern/cycles/kernel/svm/svm_ramp.h
index 563e5bcb5e4..ef8b0d103c1 100644
--- a/intern/cycles/kernel/svm/svm_ramp.h
+++ b/intern/cycles/kernel/svm/svm_ramp.h
@@ -21,13 +21,13 @@ CCL_NAMESPACE_BEGIN
 
 /* NOTE: svm_ramp.h, svm_ramp_util.h and node_ramp_util.h must stay consistent */
 
-ccl_device_inline float fetch_float(const KernelGlobals *kg, int offset)
+ccl_device_inline float fetch_float(ccl_global const KernelGlobals *kg, int offset)
 {
   uint4 node = kernel_tex_fetch(__svm_nodes, offset);
   return __uint_as_float(node.x);
 }
 
-ccl_device_inline float float_ramp_lookup(const KernelGlobals *kg,
+ccl_device_inline float float_ramp_lookup(ccl_global const KernelGlobals *kg,
                                           int offset,
                                           float f,
                                           bool interpolate,
@@ -63,7 +63,7 @@ ccl_device_inline float float_ramp_lookup(const KernelGlobals *kg,
   return a;
 }
 
-ccl_device_inline float4 rgb_ramp_lookup(const KernelGlobals *kg,
+ccl_device_inline float4 rgb_ramp_lookup(ccl_global const KernelGlobals *kg,
                                          int offset,
                                          float f,
                                          bool interpolate,
@@ -99,8 +99,11 @@ ccl_device_inline float4 rgb_ramp_lookup(const KernelGlobals *kg,
   return a;
 }
 
-ccl_device_noinline int svm_node_rgb_ramp(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_rgb_ramp(ccl_global const KernelGlobals *kg,
+                                          ccl_private ShaderData *sd,
+                                          ccl_private float *stack,
+                                          uint4 node,
+                                          int offset)
 {
   uint fac_offset, color_offset, alpha_offset;
   uint interpolate = node.z;
@@ -121,8 +124,11 @@ ccl_device_noinline int svm_node_rgb_ramp(
   return offset;
 }
 
-ccl_device_noinline int svm_node_curves(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_curves(ccl_global const KernelGlobals *kg,
+                                        ccl_private ShaderData *sd,
+                                        ccl_private float *stack,
+                                        uint4 node,
+                                        int offset)
 {
   uint fac_offset, color_offset, out_offset;
   svm_unpack_node_uchar3(node.y, &fac_offset, &color_offset, &out_offset);
@@ -147,8 +153,11 @@ ccl_device_noinline int svm_node_curves(
   return offset;
 }
 
-ccl_device_noinline int svm_node_curve(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_curve(ccl_global const KernelGlobals *kg,
+                                       ccl_private ShaderData *sd,
+                                       ccl_private float *stack,
+                                       uint4 node,
+                                       int offset)
 {
   uint fac_offset, value_in_offset, out_offset;
   svm_unpack_node_uchar3(node.y, &fac_offset, &value_in_offset, &out_offset);
diff --git a/intern/cycles/kernel/svm/svm_sepcomb_hsv.h b/intern/cycles/kernel/svm/svm_sepcomb_hsv.h
index 8d52845ea3d..3cd4ba87a55 100644
--- a/intern/cycles/kernel/svm/svm_sepcomb_hsv.h
+++ b/intern/cycles/kernel/svm/svm_sepcomb_hsv.h
@@ -16,9 +16,9 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device_noinline int svm_node_combine_hsv(const KernelGlobals *kg,
-                                             ShaderData *sd,
-                                             float *stack,
+ccl_device_noinline int svm_node_combine_hsv(ccl_global const KernelGlobals *kg,
+                                             ccl_private ShaderData *sd,
+                                             ccl_private float *stack,
                                              uint hue_in,
                                              uint saturation_in,
                                              uint value_in,
@@ -39,9 +39,9 @@ ccl_device_noinline int svm_node_combine_hsv(const KernelGlobals *kg,
   return offset;
 }
 
-ccl_device_noinline int svm_node_separate_hsv(const KernelGlobals *kg,
-                                              ShaderData *sd,
-                                              float *stack,
+ccl_device_noinline int svm_node_separate_hsv(ccl_global const KernelGlobals *kg,
+                                              ccl_private ShaderData *sd,
+                                              ccl_private float *stack,
                                               uint color_in,
                                               uint hue_out,
                                               uint saturation_out,
diff --git a/intern/cycles/kernel/svm/svm_sepcomb_vector.h b/intern/cycles/kernel/svm/svm_sepcomb_vector.h
index cbf77f1e640..11e440f2cbf 100644
--- a/intern/cycles/kernel/svm/svm_sepcomb_vector.h
+++ b/intern/cycles/kernel/svm/svm_sepcomb_vector.h
@@ -18,8 +18,11 @@ CCL_NAMESPACE_BEGIN
 
 /* Vector combine / separate, used for the RGB and XYZ nodes */
 
-ccl_device void svm_node_combine_vector(
-    ShaderData *sd, float *stack, uint in_offset, uint vector_index, uint out_offset)
+ccl_device void svm_node_combine_vector(ccl_private ShaderData *sd,
+                                        ccl_private float *stack,
+                                        uint in_offset,
+                                        uint vector_index,
+                                        uint out_offset)
 {
   float vector = stack_load_float(stack, in_offset);
 
@@ -27,8 +30,11 @@ ccl_device void svm_node_combine_vector(
     stack_store_float(stack, out_offset + vector_index, vector);
 }
 
-ccl_device void svm_node_separate_vector(
-    ShaderData *sd, float *stack, uint ivector_offset, uint vector_index, uint out_offset)
+ccl_device void svm_node_separate_vector(ccl_private ShaderData *sd,
+                                         ccl_private float *stack,
+                                         uint ivector_offset,
+                                         uint vector_index,
+                                         uint out_offset)
 {
   float3 vector = stack_load_float3(stack, ivector_offset);
 
diff --git a/intern/cycles/kernel/svm/svm_sky.h b/intern/cycles/kernel/svm/svm_sky.h
index b77c4311e72..04db8109170 100644
--- a/intern/cycles/kernel/svm/svm_sky.h
+++ b/intern/cycles/kernel/svm/svm_sky.h
@@ -28,7 +28,7 @@ ccl_device float sky_angle_between(float thetav, float phiv, float theta, float
  * "A Practical Analytic Model for Daylight"
  * A. J. Preetham, Peter Shirley, Brian Smits
  */
-ccl_device float sky_perez_function(float *lam, float theta, float gamma)
+ccl_device float sky_perez_function(ccl_private float *lam, float theta, float gamma)
 {
   float ctheta = cosf(theta);
   float cgamma = cosf(gamma);
@@ -37,16 +37,16 @@ ccl_device float sky_perez_function(float *lam, float theta, float gamma)
          (1.0f + lam[2] * expf(lam[3] * gamma) + lam[4] * cgamma * cgamma);
 }
 
-ccl_device float3 sky_radiance_preetham(const KernelGlobals *kg,
+ccl_device float3 sky_radiance_preetham(ccl_global const KernelGlobals *kg,
                                         float3 dir,
                                         float sunphi,
                                         float suntheta,
                                         float radiance_x,
                                         float radiance_y,
                                         float radiance_z,
-                                        float *config_x,
-                                        float *config_y,
-                                        float *config_z)
+                                        ccl_private float *config_x,
+                                        ccl_private float *config_y,
+                                        ccl_private float *config_z)
 {
   /* convert vector to spherical coordinates */
   float2 spherical = direction_to_spherical(dir);
@@ -73,7 +73,7 @@ ccl_device float3 sky_radiance_preetham(const KernelGlobals *kg,
  * "An Analytic Model for Full Spectral Sky-Dome Radiance"
  * Lukas Hosek, Alexander Wilkie
  */
-ccl_device float sky_radiance_internal(float *configuration, float theta, float gamma)
+ccl_device float sky_radiance_internal(ccl_private float *configuration, float theta, float gamma)
 {
   float ctheta = cosf(theta);
   float cgamma = cosf(gamma);
@@ -90,16 +90,16 @@ ccl_device float sky_radiance_internal(float *configuration, float theta, float
           configuration[6] * mieM + configuration[7] * zenith);
 }
 
-ccl_device float3 sky_radiance_hosek(const KernelGlobals *kg,
+ccl_device float3 sky_radiance_hosek(ccl_global const KernelGlobals *kg,
                                      float3 dir,
                                      float sunphi,
                                      float suntheta,
                                      float radiance_x,
                                      float radiance_y,
                                      float radiance_z,
-                                     float *config_x,
-                                     float *config_y,
-                                     float *config_z)
+                                     ccl_private float *config_x,
+                                     ccl_private float *config_y,
+                                     ccl_private float *config_z)
 {
   /* convert vector to spherical coordinates */
   float2 spherical = direction_to_spherical(dir);
@@ -127,9 +127,9 @@ ccl_device float3 geographical_to_direction(float lat, float lon)
   return make_float3(cos(lat) * cos(lon), cos(lat) * sin(lon), sin(lat));
 }
 
-ccl_device float3 sky_radiance_nishita(const KernelGlobals *kg,
+ccl_device float3 sky_radiance_nishita(ccl_global const KernelGlobals *kg,
                                        float3 dir,
-                                       float *nishita_data,
+                                       ccl_private float *nishita_data,
                                        uint texture_id)
 {
   /* definitions */
@@ -209,8 +209,11 @@ ccl_device float3 sky_radiance_nishita(const KernelGlobals *kg,
   return xyz_to_rgb(kg, xyz);
 }
 
-ccl_device_noinline int svm_node_tex_sky(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_tex_sky(ccl_global const KernelGlobals *kg,
+                                         ccl_private ShaderData *sd,
+                                         ccl_private float *stack,
+                                         uint4 node,
+                                         int offset)
 {
   /* Load data */
   uint dir_offset = node.y;
diff --git a/intern/cycles/kernel/svm/svm_tex_coord.h b/intern/cycles/kernel/svm/svm_tex_coord.h
index 8869001015b..295d5e9f65b 100644
--- a/intern/cycles/kernel/svm/svm_tex_coord.h
+++ b/intern/cycles/kernel/svm/svm_tex_coord.h
@@ -22,8 +22,12 @@ CCL_NAMESPACE_BEGIN
 
 /* Texture Coordinate Node */
 
-ccl_device_noinline int svm_node_tex_coord(
-    const KernelGlobals *kg, ShaderData *sd, int path_flag, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_tex_coord(ccl_global const KernelGlobals *kg,
+                                           ccl_private ShaderData *sd,
+                                           int path_flag,
+                                           ccl_private float *stack,
+                                           uint4 node,
+                                           int offset)
 {
   float3 data;
   uint type = node.y;
@@ -99,8 +103,12 @@ ccl_device_noinline int svm_node_tex_coord(
   return offset;
 }
 
-ccl_device_noinline int svm_node_tex_coord_bump_dx(
-    const KernelGlobals *kg, ShaderData *sd, int path_flag, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_tex_coord_bump_dx(ccl_global const KernelGlobals *kg,
+                                                   ccl_private ShaderData *sd,
+                                                   int path_flag,
+                                                   ccl_private float *stack,
+                                                   uint4 node,
+                                                   int offset)
 {
 #ifdef __RAY_DIFFERENTIALS__
   float3 data;
@@ -180,8 +188,12 @@ ccl_device_noinline int svm_node_tex_coord_bump_dx(
 #endif
 }
 
-ccl_device_noinline int svm_node_tex_coord_bump_dy(
-    const KernelGlobals *kg, ShaderData *sd, int path_flag, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_tex_coord_bump_dy(ccl_global const KernelGlobals *kg,
+                                                   ccl_private ShaderData *sd,
+                                                   int path_flag,
+                                                   ccl_private float *stack,
+                                                   uint4 node,
+                                                   int offset)
 {
 #ifdef __RAY_DIFFERENTIALS__
   float3 data;
@@ -261,9 +273,9 @@ ccl_device_noinline int svm_node_tex_coord_bump_dy(
 #endif
 }
 
-ccl_device_noinline void svm_node_normal_map(const KernelGlobals *kg,
-                                             ShaderData *sd,
-                                             float *stack,
+ccl_device_noinline void svm_node_normal_map(ccl_global const KernelGlobals *kg,
+                                             ccl_private ShaderData *sd,
+                                             ccl_private float *stack,
                                              uint4 node)
 {
   uint color_offset, strength_offset, normal_offset, space;
@@ -354,9 +366,9 @@ ccl_device_noinline void svm_node_normal_map(const KernelGlobals *kg,
   stack_store_float3(stack, normal_offset, N);
 }
 
-ccl_device_noinline void svm_node_tangent(const KernelGlobals *kg,
-                                          ShaderData *sd,
-                                          float *stack,
+ccl_device_noinline void svm_node_tangent(ccl_global const KernelGlobals *kg,
+                                          ccl_private ShaderData *sd,
+                                          ccl_private float *stack,
                                           uint4 node)
 {
   uint tangent_offset, direction_type, axis;
diff --git a/intern/cycles/kernel/svm/svm_value.h b/intern/cycles/kernel/svm/svm_value.h
index d0478660094..d1038bc072d 100644
--- a/intern/cycles/kernel/svm/svm_value.h
+++ b/intern/cycles/kernel/svm/svm_value.h
@@ -18,14 +18,20 @@ CCL_NAMESPACE_BEGIN
 
 /* Value Nodes */
 
-ccl_device void svm_node_value_f(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint ivalue, uint out_offset)
+ccl_device void svm_node_value_f(ccl_global const KernelGlobals *kg,
+                                 ccl_private ShaderData *sd,
+                                 ccl_private float *stack,
+                                 uint ivalue,
+                                 uint out_offset)
 {
   stack_store_float(stack, out_offset, __uint_as_float(ivalue));
 }
 
-ccl_device int svm_node_value_v(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint out_offset, int offset)
+ccl_device int svm_node_value_v(ccl_global const KernelGlobals *kg,
+                                ccl_private ShaderData *sd,
+                                ccl_private float *stack,
+                                uint out_offset,
+                                int offset)
 {
   /* read extra data */
   uint4 node1 = read_node(kg, &offset);
diff --git a/intern/cycles/kernel/svm/svm_vector_rotate.h b/intern/cycles/kernel/svm/svm_vector_rotate.h
index 55e1bce0158..c20f9b2556f 100644
--- a/intern/cycles/kernel/svm/svm_vector_rotate.h
+++ b/intern/cycles/kernel/svm/svm_vector_rotate.h
@@ -18,8 +18,8 @@ CCL_NAMESPACE_BEGIN
 
 /* Vector Rotate */
 
-ccl_device_noinline void svm_node_vector_rotate(ShaderData *sd,
-                                                float *stack,
+ccl_device_noinline void svm_node_vector_rotate(ccl_private ShaderData *sd,
+                                                ccl_private float *stack,
                                                 uint input_stack_offsets,
                                                 uint axis_stack_offsets,
                                                 uint result_stack_offset)
diff --git a/intern/cycles/kernel/svm/svm_vector_transform.h b/intern/cycles/kernel/svm/svm_vector_transform.h
index 8aedb7e0f54..b6c898c3952 100644
--- a/intern/cycles/kernel/svm/svm_vector_transform.h
+++ b/intern/cycles/kernel/svm/svm_vector_transform.h
@@ -18,9 +18,9 @@ CCL_NAMESPACE_BEGIN
 
 /* Vector Transform */
 
-ccl_device_noinline void svm_node_vector_transform(const KernelGlobals *kg,
-                                                   ShaderData *sd,
-                                                   float *stack,
+ccl_device_noinline void svm_node_vector_transform(ccl_global const KernelGlobals *kg,
+                                                   ccl_private ShaderData *sd,
+                                                   ccl_private float *stack,
                                                    uint4 node)
 {
   uint itype, ifrom, ito;
diff --git a/intern/cycles/kernel/svm/svm_vertex_color.h b/intern/cycles/kernel/svm/svm_vertex_color.h
index 986ea244f3a..3641f05ca43 100644
--- a/intern/cycles/kernel/svm/svm_vertex_color.h
+++ b/intern/cycles/kernel/svm/svm_vertex_color.h
@@ -16,9 +16,9 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device_noinline void svm_node_vertex_color(const KernelGlobals *kg,
-                                               ShaderData *sd,
-                                               float *stack,
+ccl_device_noinline void svm_node_vertex_color(ccl_global const KernelGlobals *kg,
+                                               ccl_private ShaderData *sd,
+                                               ccl_private float *stack,
                                                uint layer_id,
                                                uint color_offset,
                                                uint alpha_offset)
@@ -35,9 +35,9 @@ ccl_device_noinline void svm_node_vertex_color(const KernelGlobals *kg,
   }
 }
 
-ccl_device_noinline void svm_node_vertex_color_bump_dx(const KernelGlobals *kg,
-                                                       ShaderData *sd,
-                                                       float *stack,
+ccl_device_noinline void svm_node_vertex_color_bump_dx(ccl_global const KernelGlobals *kg,
+                                                       ccl_private ShaderData *sd,
+                                                       ccl_private float *stack,
                                                        uint layer_id,
                                                        uint color_offset,
                                                        uint alpha_offset)
@@ -56,9 +56,9 @@ ccl_device_noinline void svm_node_vertex_color_bump_dx(const KernelGlobals *kg,
   }
 }
 
-ccl_device_noinline void svm_node_vertex_color_bump_dy(const KernelGlobals *kg,
-                                                       ShaderData *sd,
-                                                       float *stack,
+ccl_device_noinline void svm_node_vertex_color_bump_dy(ccl_global const KernelGlobals *kg,
+                                                       ccl_private ShaderData *sd,
+                                                       ccl_private float *stack,
                                                        uint layer_id,
                                                        uint color_offset,
                                                        uint alpha_offset)
diff --git a/intern/cycles/kernel/svm/svm_voronoi.h b/intern/cycles/kernel/svm/svm_voronoi.h
index b1d2eff7f37..e7112087e17 100644
--- a/intern/cycles/kernel/svm/svm_voronoi.h
+++ b/intern/cycles/kernel/svm/svm_voronoi.h
@@ -46,9 +46,9 @@ ccl_device void voronoi_f1_1d(float w,
                               float exponent,
                               float randomness,
                               NodeVoronoiDistanceMetric metric,
-                              float *outDistance,
-                              float3 *outColor,
-                              float *outW)
+                              ccl_private float *outDistance,
+                              ccl_private float3 *outColor,
+                              ccl_private float *outW)
 {
   float cellPosition = floorf(w);
   float localPosition = w - cellPosition;
@@ -76,9 +76,9 @@ ccl_device void voronoi_smooth_f1_1d(float w,
                                      float exponent,
                                      float randomness,
                                      NodeVoronoiDistanceMetric metric,
-                                     float *outDistance,
-                                     float3 *outColor,
-                                     float *outW)
+                                     ccl_private float *outDistance,
+                                     ccl_private float3 *outColor,
+                                     ccl_private float *outW)
 {
   float cellPosition = floorf(w);
   float localPosition = w - cellPosition;
@@ -108,9 +108,9 @@ ccl_device void voronoi_f2_1d(float w,
                               float exponent,
                               float randomness,
                               NodeVoronoiDistanceMetric metric,
-                              float *outDistance,
-                              float3 *outColor,
-                              float *outW)
+                              ccl_private float *outDistance,
+                              ccl_private float3 *outColor,
+                              ccl_private float *outW)
 {
   float cellPosition = floorf(w);
   float localPosition = w - cellPosition;
@@ -144,7 +144,9 @@ ccl_device void voronoi_f2_1d(float w,
   *outW = positionF2 + cellPosition;
 }
 
-ccl_device void voronoi_distance_to_edge_1d(float w, float randomness, float *outDistance)
+ccl_device void voronoi_distance_to_edge_1d(float w,
+                                            float randomness,
+                                            ccl_private float *outDistance)
 {
   float cellPosition = floorf(w);
   float localPosition = w - cellPosition;
@@ -158,7 +160,7 @@ ccl_device void voronoi_distance_to_edge_1d(float w, float randomness, float *ou
   *outDistance = min(distanceToMidLeft, distanceToMidRight);
 }
 
-ccl_device void voronoi_n_sphere_radius_1d(float w, float randomness, float *outRadius)
+ccl_device void voronoi_n_sphere_radius_1d(float w, float randomness, ccl_private float *outRadius)
 {
   float cellPosition = floorf(w);
   float localPosition = w - cellPosition;
@@ -223,9 +225,9 @@ ccl_device void voronoi_f1_2d(float2 coord,
                               float exponent,
                               float randomness,
                               NodeVoronoiDistanceMetric metric,
-                              float *outDistance,
-                              float3 *outColor,
-                              float2 *outPosition)
+                              ccl_private float *outDistance,
+                              ccl_private float3 *outColor,
+                              ccl_private float2 *outPosition)
 {
   float2 cellPosition = floor(coord);
   float2 localPosition = coord - cellPosition;
@@ -256,9 +258,9 @@ ccl_device void voronoi_smooth_f1_2d(float2 coord,
                                      float exponent,
                                      float randomness,
                                      NodeVoronoiDistanceMetric metric,
-                                     float *outDistance,
-                                     float3 *outColor,
-                                     float2 *outPosition)
+                                     ccl_private float *outDistance,
+                                     ccl_private float3 *outColor,
+                                     ccl_private float2 *outPosition)
 {
   float2 cellPosition = floor(coord);
   float2 localPosition = coord - cellPosition;
@@ -291,9 +293,9 @@ ccl_device void voronoi_f2_2d(float2 coord,
                               float exponent,
                               float randomness,
                               NodeVoronoiDistanceMetric metric,
-                              float *outDistance,
-                              float3 *outColor,
-                              float2 *outPosition)
+                              ccl_private float *outDistance,
+                              ccl_private float3 *outColor,
+                              ccl_private float2 *outPosition)
 {
   float2 cellPosition = floor(coord);
   float2 localPosition = coord - cellPosition;
@@ -330,7 +332,9 @@ ccl_device void voronoi_f2_2d(float2 coord,
   *outPosition = positionF2 + cellPosition;
 }
 
-ccl_device void voronoi_distance_to_edge_2d(float2 coord, float randomness, float *outDistance)
+ccl_device void voronoi_distance_to_edge_2d(float2 coord,
+                                            float randomness,
+                                            ccl_private float *outDistance)
 {
   float2 cellPosition = floor(coord);
   float2 localPosition = coord - cellPosition;
@@ -369,7 +373,9 @@ ccl_device void voronoi_distance_to_edge_2d(float2 coord, float randomness, floa
   *outDistance = minDistance;
 }
 
-ccl_device void voronoi_n_sphere_radius_2d(float2 coord, float randomness, float *outRadius)
+ccl_device void voronoi_n_sphere_radius_2d(float2 coord,
+                                           float randomness,
+                                           ccl_private float *outRadius)
 {
   float2 cellPosition = floor(coord);
   float2 localPosition = coord - cellPosition;
@@ -441,9 +447,9 @@ ccl_device void voronoi_f1_3d(float3 coord,
                               float exponent,
                               float randomness,
                               NodeVoronoiDistanceMetric metric,
-                              float *outDistance,
-                              float3 *outColor,
-                              float3 *outPosition)
+                              ccl_private float *outDistance,
+                              ccl_private float3 *outColor,
+                              ccl_private float3 *outPosition)
 {
   float3 cellPosition = floor(coord);
   float3 localPosition = coord - cellPosition;
@@ -477,9 +483,9 @@ ccl_device void voronoi_smooth_f1_3d(float3 coord,
                                      float exponent,
                                      float randomness,
                                      NodeVoronoiDistanceMetric metric,
-                                     float *outDistance,
-                                     float3 *outColor,
-                                     float3 *outPosition)
+                                     ccl_private float *outDistance,
+                                     ccl_private float3 *outColor,
+                                     ccl_private float3 *outPosition)
 {
   float3 cellPosition = floor(coord);
   float3 localPosition = coord - cellPosition;
@@ -515,9 +521,9 @@ ccl_device void voronoi_f2_3d(float3 coord,
                               float exponent,
                               float randomness,
                               NodeVoronoiDistanceMetric metric,
-                              float *outDistance,
-                              float3 *outColor,
-                              float3 *outPosition)
+                              ccl_private float *outDistance,
+                              ccl_private float3 *outColor,
+                              ccl_private float3 *outPosition)
 {
   float3 cellPosition = floor(coord);
   float3 localPosition = coord - cellPosition;
@@ -557,7 +563,9 @@ ccl_device void voronoi_f2_3d(float3 coord,
   *outPosition = positionF2 + cellPosition;
 }
 
-ccl_device void voronoi_distance_to_edge_3d(float3 coord, float randomness, float *outDistance)
+ccl_device void voronoi_distance_to_edge_3d(float3 coord,
+                                            float randomness,
+                                            ccl_private float *outDistance)
 {
   float3 cellPosition = floor(coord);
   float3 localPosition = coord - cellPosition;
@@ -600,7 +608,9 @@ ccl_device void voronoi_distance_to_edge_3d(float3 coord, float randomness, floa
   *outDistance = minDistance;
 }
 
-ccl_device void voronoi_n_sphere_radius_3d(float3 coord, float randomness, float *outRadius)
+ccl_device void voronoi_n_sphere_radius_3d(float3 coord,
+                                           float randomness,
+                                           ccl_private float *outRadius)
 {
   float3 cellPosition = floor(coord);
   float3 localPosition = coord - cellPosition;
@@ -676,9 +686,9 @@ ccl_device void voronoi_f1_4d(float4 coord,
                               float exponent,
                               float randomness,
                               NodeVoronoiDistanceMetric metric,
-                              float *outDistance,
-                              float3 *outColor,
-                              float4 *outPosition)
+                              ccl_private float *outDistance,
+                              ccl_private float3 *outColor,
+                              ccl_private float4 *outPosition)
 {
   float4 cellPosition = floor(coord);
   float4 localPosition = coord - cellPosition;
@@ -715,9 +725,9 @@ ccl_device void voronoi_smooth_f1_4d(float4 coord,
                                      float exponent,
                                      float randomness,
                                      NodeVoronoiDistanceMetric metric,
-                                     float *outDistance,
-                                     float3 *outColor,
-                                     float4 *outPosition)
+                                     ccl_private float *outDistance,
+                                     ccl_private float3 *outColor,
+                                     ccl_private float4 *outPosition)
 {
   float4 cellPosition = floor(coord);
   float4 localPosition = coord - cellPosition;
@@ -756,9 +766,9 @@ ccl_device void voronoi_f2_4d(float4 coord,
                               float exponent,
                               float randomness,
                               NodeVoronoiDistanceMetric metric,
-                              float *outDistance,
-                              float3 *outColor,
-                              float4 *outPosition)
+                              ccl_private float *outDistance,
+                              ccl_private float3 *outColor,
+                              ccl_private float4 *outPosition)
 {
   float4 cellPosition = floor(coord);
   float4 localPosition = coord - cellPosition;
@@ -801,7 +811,9 @@ ccl_device void voronoi_f2_4d(float4 coord,
   *outPosition = positionF2 + cellPosition;
 }
 
-ccl_device void voronoi_distance_to_edge_4d(float4 coord, float randomness, float *outDistance)
+ccl_device void voronoi_distance_to_edge_4d(float4 coord,
+                                            float randomness,
+                                            ccl_private float *outDistance)
 {
   float4 cellPosition = floor(coord);
   float4 localPosition = coord - cellPosition;
@@ -850,7 +862,9 @@ ccl_device void voronoi_distance_to_edge_4d(float4 coord, float randomness, floa
   *outDistance = minDistance;
 }
 
-ccl_device void voronoi_n_sphere_radius_4d(float4 coord, float randomness, float *outRadius)
+ccl_device void voronoi_n_sphere_radius_4d(float4 coord,
+                                           float randomness,
+                                           ccl_private float *outRadius)
 {
   float4 cellPosition = floor(coord);
   float4 localPosition = coord - cellPosition;
@@ -903,9 +917,9 @@ ccl_device void voronoi_n_sphere_radius_4d(float4 coord, float randomness, float
 }
 
 template<uint node_feature_mask>
-ccl_device_noinline int svm_node_tex_voronoi(const KernelGlobals *kg,
-                                             ShaderData *sd,
-                                             float *stack,
+ccl_device_noinline int svm_node_tex_voronoi(ccl_global const KernelGlobals *kg,
+                                             ccl_private ShaderData *sd,
+                                             ccl_private float *stack,
                                              uint dimensions,
                                              uint feature,
                                              uint metric,
diff --git a/intern/cycles/kernel/svm/svm_voxel.h b/intern/cycles/kernel/svm/svm_voxel.h
index 78b75405356..764fb71ba72 100644
--- a/intern/cycles/kernel/svm/svm_voxel.h
+++ b/intern/cycles/kernel/svm/svm_voxel.h
@@ -19,8 +19,11 @@ CCL_NAMESPACE_BEGIN
 /* TODO(sergey): Think of making it more generic volume-type attribute
  * sampler.
  */
-ccl_device_noinline int svm_node_tex_voxel(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_tex_voxel(ccl_global const KernelGlobals *kg,
+                                           ccl_private ShaderData *sd,
+                                           ccl_private float *stack,
+                                           uint4 node,
+                                           int offset)
 {
   uint co_offset, density_out_offset, color_out_offset, space;
   svm_unpack_node_uchar4(node.z, &co_offset, &density_out_offset, &color_out_offset, &space);
diff --git a/intern/cycles/kernel/svm/svm_wave.h b/intern/cycles/kernel/svm/svm_wave.h
index 00f980c16df..1ac130e2006 100644
--- a/intern/cycles/kernel/svm/svm_wave.h
+++ b/intern/cycles/kernel/svm/svm_wave.h
@@ -82,8 +82,11 @@ ccl_device_noinline_cpu float svm_wave(NodeWaveType type,
   }
 }
 
-ccl_device_noinline int svm_node_tex_wave(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int offset)
+ccl_device_noinline int svm_node_tex_wave(ccl_global const KernelGlobals *kg,
+                                          ccl_private ShaderData *sd,
+                                          ccl_private float *stack,
+                                          uint4 node,
+                                          int offset)
 {
   uint4 node2 = read_node(kg, &offset);
   uint4 node3 = read_node(kg, &offset);
diff --git a/intern/cycles/kernel/svm/svm_wavelength.h b/intern/cycles/kernel/svm/svm_wavelength.h
index aa291fd2741..e891744f276 100644
--- a/intern/cycles/kernel/svm/svm_wavelength.h
+++ b/intern/cycles/kernel/svm/svm_wavelength.h
@@ -34,8 +34,11 @@ CCL_NAMESPACE_BEGIN
 
 /* Wavelength to RGB */
 
-ccl_device_noinline void svm_node_wavelength(
-    const KernelGlobals *kg, ShaderData *sd, float *stack, uint wavelength, uint color_out)
+ccl_device_noinline void svm_node_wavelength(ccl_global const KernelGlobals *kg,
+                                             ccl_private ShaderData *sd,
+                                             ccl_private float *stack,
+                                             uint wavelength,
+                                             uint color_out)
 {
   // CIE colour matching functions xBar, yBar, and zBar for
   //   wavelengths from 380 through 780 nanometers, every 5
diff --git a/intern/cycles/kernel/svm/svm_white_noise.h b/intern/cycles/kernel/svm/svm_white_noise.h
index 0306d2e7b9c..ccc49bf1a7c 100644
--- a/intern/cycles/kernel/svm/svm_white_noise.h
+++ b/intern/cycles/kernel/svm/svm_white_noise.h
@@ -16,9 +16,9 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device_noinline void svm_node_tex_white_noise(const KernelGlobals *kg,
-                                                  ShaderData *sd,
-                                                  float *stack,
+ccl_device_noinline void svm_node_tex_white_noise(ccl_global const KernelGlobals *kg,
+                                                  ccl_private ShaderData *sd,
+                                                  ccl_private float *stack,
                                                   uint dimensions,
                                                   uint inputs_stack_offsets,
                                                   uint ouptuts_stack_offsets)
diff --git a/intern/cycles/kernel/svm/svm_wireframe.h b/intern/cycles/kernel/svm/svm_wireframe.h
index 7ec913789d2..70d1211aa4a 100644
--- a/intern/cycles/kernel/svm/svm_wireframe.h
+++ b/intern/cycles/kernel/svm/svm_wireframe.h
@@ -34,8 +34,11 @@ CCL_NAMESPACE_BEGIN
 
 /* Wireframe Node */
 
-ccl_device_inline float wireframe(
-    const KernelGlobals *kg, ShaderData *sd, float size, int pixel_size, float3 *P)
+ccl_device_inline float wireframe(ccl_global const KernelGlobals *kg,
+                                  ccl_private ShaderData *sd,
+                                  float size,
+                                  int pixel_size,
+                                  ccl_private float3 *P)
 {
 #ifdef __HAIR__
   if (sd->prim != PRIM_NONE && sd->type & PRIMITIVE_ALL_TRIANGLE)
@@ -88,9 +91,9 @@ ccl_device_inline float wireframe(
   return 0.0f;
 }
 
-ccl_device_noinline void svm_node_wireframe(const KernelGlobals *kg,
-                                            ShaderData *sd,
-                                            float *stack,
+ccl_device_noinline void svm_node_wireframe(ccl_global const KernelGlobals *kg,
+                                            ccl_private ShaderData *sd,
+                                            ccl_private float *stack,
                                             uint4 node)
 {
   uint in_size = node.y;