Merge branch 'draw-viewport-data' into eevee-rewrite

# Conflicts:
#	release/scripts/startup/bl_ui/properties_data_camera.py
#	source/blender/blenkernel/BKE_camera.h
#	source/blender/blenkernel/BKE_node.h
#	source/blender/blenkernel/intern/camera.c
#	source/blender/blenlib/BLI_float2.hh
#	source/blender/blenlib/BLI_float3.hh
#	source/blender/blenlib/BLI_float4.hh
#	source/blender/blenlib/BLI_math_geom.h
#	source/blender/blenlib/intern/math_geom.c
#	source/blender/draw/CMakeLists.txt
#	source/blender/draw/engines/basic/basic_engine.c
#	source/blender/draw/engines/eevee/eevee_cryptomatte.c
#	source/blender/draw/engines/eevee/eevee_effects.c
#	source/blender/draw/engines/eevee/eevee_engine.c
#	source/blender/draw/engines/eevee/eevee_lightcache.c
#	source/blender/draw/engines/eevee/eevee_lightcache.h
#	source/blender/draw/engines/eevee/eevee_lightprobes.c
#	source/blender/draw/engines/eevee/eevee_lights.c
#	source/blender/draw/engines/eevee/eevee_materials.c
#	source/blender/draw/engines/eevee/eevee_motion_blur.c
#	source/blender/draw/engines/eevee/eevee_occlusion.c
#	source/blender/draw/engines/eevee/eevee_private.h
#	source/blender/draw/engines/eevee/eevee_render.c
#	source/blender/draw/engines/eevee/eevee_renderpasses.c
#	source/blender/draw/engines/eevee/eevee_sampling.c
#	source/blender/draw/engines/eevee/eevee_screen_raytrace.c
#	source/blender/draw/engines/eevee/eevee_shaders.c
#	source/blender/draw/engines/eevee/eevee_shadows.c
#	source/blender/draw/engines/eevee/eevee_shadows_cube.c
#	source/blender/draw/engines/eevee/eevee_temporal_sampling.c
#	source/blender/draw/engines/eevee/shaders/ambient_occlusion_lib.glsl
#	source/blender/draw/engines/eevee/shaders/closure_eval_lib.glsl
#	source/blender/draw/engines/eevee/shaders/common_utiltex_lib.glsl
#	source/blender/draw/engines/eevee/shaders/effect_dof_bokeh_frag.glsl
#	source/blender/draw/engines/eevee/shaders/effect_dof_gather_frag.glsl
#	source/blender/draw/engines/eevee/shaders/effect_dof_reduce_frag.glsl
#	source/blender/draw/engines/eevee/shaders/effect_reflection_resolve_frag.glsl
#	source/blender/draw/engines/eevee/shaders/effect_temporal_aa.glsl
#	source/blender/draw/engines/eevee/shaders/random_lib.glsl
#	source/blender/draw/engines/eevee/shaders/shadow_vert.glsl
#	source/blender/draw/engines/eevee/shaders/surface_lib.glsl
#	source/blender/draw/engines/eevee/shaders/surface_vert.glsl
#	source/blender/draw/engines/eevee/shaders/volumetric_lib.glsl
#	source/blender/draw/engines/external/external_engine.c
#	source/blender/draw/engines/gpencil/gpencil_engine.c
#	source/blender/draw/engines/image/image_engine.c
#	source/blender/draw/engines/overlay/overlay_engine.c
#	source/blender/draw/engines/select/select_engine.c
#	source/blender/draw/engines/workbench/shaders/workbench_volume_frag.glsl
#	source/blender/draw/engines/workbench/shaders/workbench_volume_vert.glsl
#	source/blender/draw/engines/workbench/workbench_engine.c
#	source/blender/draw/engines/workbench/workbench_shader.c
#	source/blender/draw/intern/DRW_render.h
#	source/blender/draw/intern/draw_debug.h
#	source/blender/draw/intern/draw_manager_data.c
#	source/blender/draw/intern/draw_manager_exec.c
#	source/blender/draw/intern/draw_view_data.h
#	source/blender/gpu/CMakeLists.txt
#	source/blender/gpu/GPU_material.h
#	source/blender/gpu/GPU_shader.h
#	source/blender/gpu/GPU_state.h
#	source/blender/gpu/GPU_vertex_buffer.h
#	source/blender/gpu/intern/gpu_codegen.c
#	source/blender/gpu/intern/gpu_material.c
#	source/blender/gpu/intern/gpu_material_library.h
#	source/blender/gpu/intern/gpu_node_graph.c
#	source/blender/gpu/intern/gpu_texture_private.hh
#	source/blender/gpu/intern/gpu_vertex_buffer.cc
#	source/blender/gpu/opengl/gl_shader.cc
#	source/blender/gpu/shaders/gpu_shader_common_obinfos_lib.glsl
#	source/blender/gpu/shaders/material/gpu_shader_material_shader_to_rgba.glsl
#	source/blender/nodes/shader/node_shader_tree.cc
#	source/blender/nodes/shader/nodes/node_shader_background.cc
#	source/blender/nodes/shader/nodes/node_shader_bsdf_anisotropic.cc
#	source/blender/nodes/shader/nodes/node_shader_bsdf_diffuse.cc
#	source/blender/nodes/shader/nodes/node_shader_bsdf_glass.cc
#	source/blender/nodes/shader/nodes/node_shader_bsdf_glossy.cc
#	source/blender/nodes/shader/nodes/node_shader_bsdf_hair.cc
#	source/blender/nodes/shader/nodes/node_shader_bsdf_hair_principled.c
#	source/blender/nodes/shader/nodes/node_shader_bsdf_principled.c
#	source/blender/nodes/shader/nodes/node_shader_bsdf_refraction.cc
#	source/blender/nodes/shader/nodes/node_shader_bsdf_toon.cc
#	source/blender/nodes/shader/nodes/node_shader_bsdf_translucent.cc
#	source/blender/nodes/shader/nodes/node_shader_bsdf_transparent.cc
#	source/blender/nodes/shader/nodes/node_shader_bsdf_velvet.cc
#	source/blender/nodes/shader/nodes/node_shader_eevee_specular.cc
#	source/blender/nodes/shader/nodes/node_shader_emission.cc
#	source/blender/nodes/shader/nodes/node_shader_holdout.cc
#	source/blender/nodes/shader/nodes/node_shader_output_material.cc
#	source/blender/nodes/shader/nodes/node_shader_subsurface_scattering.c
#	source/blender/nodes/shader/nodes/node_shader_tex_coord.cc
#	source/blender/nodes/shader/nodes/node_shader_vector_transform.cc
#	source/blender/nodes/shader/nodes/node_shader_volume_absorption.cc
#	source/blender/nodes/shader/nodes/node_shader_volume_principled.cc
#	source/blender/nodes/shader/nodes/node_shader_volume_scatter.cc
#	source/blender/render/RE_pipeline.h
#	source/blender/render/intern/initrender.c

1 files changed, 9 insertions, 5 deletions

diff --git a/intern/cycles/kernel/device/gpu/parallel_prefix_sum.h b/intern/cycles/kernel/device/gpu/parallel_prefix_sum.h
index a1349e82efb..4bd002c27e4 100644
--- a/intern/cycles/kernel/device/gpu/parallel_prefix_sum.h
+++ b/intern/cycles/kernel/device/gpu/parallel_prefix_sum.h
@@ -25,7 +25,7 @@ CCL_NAMESPACE_BEGIN
  * This is used for an array the size of the number of shaders in the scene
  * which is not usually huge, so might not be a significant bottleneck. */
 
-#include "util/util_atomic.h"
+#include "util/atomic.h"
 
 #ifdef __HIP__
 #  define GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE 1024
@@ -33,16 +33,20 @@ CCL_NAMESPACE_BEGIN
 #  define GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE 512
 #endif
 
-template<uint blocksize> __device__ void gpu_parallel_prefix_sum(int *values, const int num_values)
+__device__ void gpu_parallel_prefix_sum(const int global_id,
+                                        ccl_global int *counter,
+                                        ccl_global int *prefix_sum,
+                                        const int num_values)
 {
-  if (!(ccl_gpu_block_idx_x == 0 && ccl_gpu_thread_idx_x == 0)) {
+  if (global_id != 0) {
     return;
   }
 
   int offset = 0;
   for (int i = 0; i < num_values; i++) {
-    const int new_offset = offset + values[i];
-    values[i] = offset;
+    const int new_offset = offset + counter[i];
+    prefix_sum[i] = offset;
+    counter[i] = 0;
     offset = new_offset;
   }
 }