Blender  V2.93
kernel_id_passes.h
Go to the documentation of this file.
1 /*
2  * Copyright 2018 Blender Foundation
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16 
18 
20  int num_slots,
21  float id,
22  float weight)
23 {
24  kernel_assert(id != ID_NONE);
25  if (weight == 0.0f) {
26  return;
27  }
28 
29  for (int slot = 0; slot < num_slots; slot++) {
30  ccl_global float2 *id_buffer = (ccl_global float2 *)buffer;
31 #ifdef __ATOMIC_PASS_WRITE__
32  /* If the loop reaches an empty slot, the ID isn't in any slot yet - so add it! */
33  if (id_buffer[slot].x == ID_NONE) {
34  /* Use an atomic to claim this slot.
35  * If a different thread got here first, try again from this slot on. */
36  float old_id = atomic_compare_and_swap_float(buffer + slot * 2, ID_NONE, id);
37  if (old_id != ID_NONE && old_id != id) {
38  continue;
39  }
40  atomic_add_and_fetch_float(buffer + slot * 2 + 1, weight);
41  break;
42  }
43  /* If there already is a slot for that ID, add the weight.
44  * If no slot was found, add it to the last. */
45  else if (id_buffer[slot].x == id || slot == num_slots - 1) {
46  atomic_add_and_fetch_float(buffer + slot * 2 + 1, weight);
47  break;
48  }
49 #else /* __ATOMIC_PASS_WRITE__ */
50  /* If the loop reaches an empty slot, the ID isn't in any slot yet - so add it! */
51  if (id_buffer[slot].x == ID_NONE) {
52  id_buffer[slot].x = id;
53  id_buffer[slot].y = weight;
54  break;
55  }
56  /* If there already is a slot for that ID, add the weight.
57  * If no slot was found, add it to the last. */
58  else if (id_buffer[slot].x == id || slot == num_slots - 1) {
59  id_buffer[slot].y += weight;
60  break;
61  }
62 #endif /* __ATOMIC_PASS_WRITE__ */
63  }
64 }
65 
67 {
68  ccl_global float2 *id_buffer = (ccl_global float2 *)buffer;
69  for (int slot = 1; slot < num_slots; ++slot) {
70  if (id_buffer[slot].x == ID_NONE) {
71  return;
72  }
73  /* Since we're dealing with a tiny number of elements, insertion sort should be fine. */
74  int i = slot;
75  while (i > 0 && id_buffer[i].y > id_buffer[i - 1].y) {
76  float2 swap = id_buffer[i];
77  id_buffer[i] = id_buffer[i - 1];
78  id_buffer[i - 1] = swap;
79  --i;
80  }
81  }
82 }
83 
84 #ifdef __KERNEL_GPU__
85 /* post-sorting for Cryptomatte */
86 ccl_device void kernel_cryptomatte_post(
87  KernelGlobals *kg, ccl_global float *buffer, uint sample, int x, int y, int offset, int stride)
88 {
89  if (sample - 1 == kernel_data.integrator.aa_samples) {
90  int index = offset + x + y * stride;
91  int pass_stride = kernel_data.film.pass_stride;
92  ccl_global float *cryptomatte_buffer = buffer + index * pass_stride +
93  kernel_data.film.pass_cryptomatte;
94  kernel_sort_id_slots(cryptomatte_buffer, 2 * kernel_data.film.cryptomatte_depth);
95  }
96 }
97 #endif
98 
unsigned int uint
Definition: BLI_sys_types.h:83
void swap(T &a, T &b)
Definition: Common.h:33
_GL_VOID GLfloat value _GL_VOID_RET _GL_VOID const GLuint GLboolean *residences _GL_BOOL_RET _GL_VOID GLsizei GLfloat GLfloat GLfloat GLfloat const GLubyte *bitmap _GL_VOID_RET _GL_VOID GLenum const void *lists _GL_VOID_RET _GL_VOID const GLdouble *equation _GL_VOID_RET _GL_VOID GLdouble GLdouble blue _GL_VOID_RET _GL_VOID GLfloat GLfloat blue _GL_VOID_RET _GL_VOID GLint GLint blue _GL_VOID_RET _GL_VOID GLshort GLshort blue _GL_VOID_RET _GL_VOID GLubyte GLubyte blue _GL_VOID_RET _GL_VOID GLuint GLuint blue _GL_VOID_RET _GL_VOID GLushort GLushort blue _GL_VOID_RET _GL_VOID GLbyte GLbyte GLbyte alpha _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble alpha _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat alpha _GL_VOID_RET _GL_VOID GLint GLint GLint alpha _GL_VOID_RET _GL_VOID GLshort GLshort GLshort alpha _GL_VOID_RET _GL_VOID GLubyte GLubyte GLubyte alpha _GL_VOID_RET _GL_VOID GLuint GLuint GLuint alpha _GL_VOID_RET _GL_VOID GLushort GLushort GLushort alpha _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLint y
_GL_VOID GLfloat value _GL_VOID_RET _GL_VOID const GLuint GLboolean *residences _GL_BOOL_RET _GL_VOID GLsizei GLfloat GLfloat GLfloat GLfloat const GLubyte *bitmap _GL_VOID_RET _GL_VOID GLenum const void *lists _GL_VOID_RET _GL_VOID const GLdouble *equation _GL_VOID_RET _GL_VOID GLdouble GLdouble blue _GL_VOID_RET _GL_VOID GLfloat GLfloat blue _GL_VOID_RET _GL_VOID GLint GLint blue _GL_VOID_RET _GL_VOID GLshort GLshort blue _GL_VOID_RET _GL_VOID GLubyte GLubyte blue _GL_VOID_RET _GL_VOID GLuint GLuint blue _GL_VOID_RET _GL_VOID GLushort GLushort blue _GL_VOID_RET _GL_VOID GLbyte GLbyte GLbyte alpha _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble alpha _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat alpha _GL_VOID_RET _GL_VOID GLint GLint GLint alpha _GL_VOID_RET _GL_VOID GLshort GLshort GLshort alpha _GL_VOID_RET _GL_VOID GLubyte GLubyte GLubyte alpha _GL_VOID_RET _GL_VOID GLuint GLuint GLuint alpha _GL_VOID_RET _GL_VOID GLushort GLushort GLushort alpha _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLint GLsizei GLsizei GLenum type _GL_VOID_RET _GL_VOID GLsizei GLenum GLenum const void *pixels _GL_VOID_RET _GL_VOID const void *pointer _GL_VOID_RET _GL_VOID GLdouble v _GL_VOID_RET _GL_VOID GLfloat v _GL_VOID_RET _GL_VOID GLint GLint i2 _GL_VOID_RET _GL_VOID GLint j _GL_VOID_RET _GL_VOID GLfloat param _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble GLdouble GLdouble zFar _GL_VOID_RET _GL_UINT GLdouble *equation _GL_VOID_RET _GL_VOID GLenum GLint *params _GL_VOID_RET _GL_VOID GLenum GLfloat *v _GL_VOID_RET _GL_VOID GLenum GLfloat *params _GL_VOID_RET _GL_VOID GLfloat *values _GL_VOID_RET _GL_VOID GLushort *values _GL_VOID_RET _GL_VOID GLenum GLfloat *params _GL_VOID_RET _GL_VOID GLenum GLdouble *params _GL_VOID_RET _GL_VOID GLenum GLint *params _GL_VOID_RET _GL_VOID GLsizei stride
#define kernel_data
#define kernel_assert(cond)
#define ccl_device
#define ccl_device_inline
#define ccl_global
#define CCL_NAMESPACE_END
ccl_device_inline void kernel_sort_id_slots(ccl_global float *buffer, int num_slots)
CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_write_id_slots(ccl_global float *buffer, int num_slots, float id, float weight)
__kernel void ccl_constant KernelData ccl_global void ccl_global char ccl_global int ccl_global char ccl_global unsigned int ccl_global float * buffer
#define ID_NONE
Definition: kernel_types.h:62
static void sample(SocketReader *reader, int x, int y, float color[4])
#define atomic_compare_and_swap_float(p, old_val, new_val)
Definition: util_atomic.h:26
#define atomic_add_and_fetch_float(p, x)
Definition: util_atomic.h:25