Blender  V2.93
COM_VariableSizeBokehBlurOperation.cc
Go to the documentation of this file.
1 /*
2  * This program is free software; you can redistribute it and/or
3  * modify it under the terms of the GNU General Public License
4  * as published by the Free Software Foundation; either version 2
5  * of the License, or (at your option) any later version.
6  *
7  * This program is distributed in the hope that it will be useful,
8  * but WITHOUT ANY WARRANTY; without even the implied warranty of
9  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
10  * GNU General Public License for more details.
11  *
12  * You should have received a copy of the GNU General Public License
13  * along with this program; if not, write to the Free Software Foundation,
14  * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
15  *
16  * Copyright 2011, Blender Foundation.
17  */
18 
20 #include "BLI_math.h"
21 #include "COM_OpenCLDevice.h"
22 
23 #include "RE_pipeline.h"
24 
25 namespace blender::compositor {
26 
28 {
30  this->addInputSocket(DataType::Color, ResizeMode::None); // do not resize the bokeh image.
31  this->addInputSocket(DataType::Value); // radius
32 #ifdef COM_DEFOCUS_SEARCH
34  ResizeMode::None); // inverse search radius optimization structure.
35 #endif
37  flags.complex = true;
38  flags.open_cl = true;
39 
40  this->m_inputProgram = nullptr;
41  this->m_inputBokehProgram = nullptr;
42  this->m_inputSizeProgram = nullptr;
43  this->m_maxBlur = 32.0f;
44  this->m_threshold = 1.0f;
45  this->m_do_size_scale = false;
46 #ifdef COM_DEFOCUS_SEARCH
47  this->m_inputSearchProgram = nullptr;
48 #endif
49 }
50 
52 {
53  this->m_inputProgram = getInputSocketReader(0);
54  this->m_inputBokehProgram = getInputSocketReader(1);
55  this->m_inputSizeProgram = getInputSocketReader(2);
56 #ifdef COM_DEFOCUS_SEARCH
57  this->m_inputSearchProgram = getInputSocketReader(3);
58 #endif
60 }
66 };
67 
69 {
71  data->color = (MemoryBuffer *)this->m_inputProgram->initializeTileData(rect);
72  data->bokeh = (MemoryBuffer *)this->m_inputBokehProgram->initializeTileData(rect);
73  data->size = (MemoryBuffer *)this->m_inputSizeProgram->initializeTileData(rect);
74 
75  rcti rect2;
77  rect, (ReadBufferOperation *)this->m_inputSizeProgram, &rect2);
78 
79  const float max_dim = MAX2(m_width, m_height);
80  const float scalar = this->m_do_size_scale ? (max_dim / 100.0f) : 1.0f;
81 
82  data->maxBlurScalar = (int)(data->size->get_max_value(rect2) * scalar);
83  CLAMP(data->maxBlurScalar, 1.0f, this->m_maxBlur);
84  return data;
85 }
86 
88 {
90  delete result;
91 }
92 
94 {
96  MemoryBuffer *inputProgramBuffer = tileData->color;
97  MemoryBuffer *inputBokehBuffer = tileData->bokeh;
98  MemoryBuffer *inputSizeBuffer = tileData->size;
99  float *inputSizeFloatBuffer = inputSizeBuffer->getBuffer();
100  float *inputProgramFloatBuffer = inputProgramBuffer->getBuffer();
101  float readColor[4];
102  float bokeh[4];
103  float tempSize[4];
104  float multiplier_accum[4];
105  float color_accum[4];
106 
107  const float max_dim = MAX2(m_width, m_height);
108  const float scalar = this->m_do_size_scale ? (max_dim / 100.0f) : 1.0f;
109  int maxBlurScalar = tileData->maxBlurScalar;
110 
111  BLI_assert(inputBokehBuffer->getWidth() == COM_BLUR_BOKEH_PIXELS);
112  BLI_assert(inputBokehBuffer->getHeight() == COM_BLUR_BOKEH_PIXELS);
113 
114 #ifdef COM_DEFOCUS_SEARCH
115  float search[4];
116  this->m_inputSearchProgram->read(search,
117  x / InverseSearchRadiusOperation::DIVIDER,
118  y / InverseSearchRadiusOperation::DIVIDER,
119  nullptr);
120  int minx = search[0];
121  int miny = search[1];
122  int maxx = search[2];
123  int maxy = search[3];
124 #else
125  int minx = MAX2(x - maxBlurScalar, 0);
126  int miny = MAX2(y - maxBlurScalar, 0);
127  int maxx = MIN2(x + maxBlurScalar, (int)m_width);
128  int maxy = MIN2(y + maxBlurScalar, (int)m_height);
129 #endif
130  {
131  inputSizeBuffer->readNoCheck(tempSize, x, y);
132  inputProgramBuffer->readNoCheck(readColor, x, y);
133 
134  copy_v4_v4(color_accum, readColor);
135  copy_v4_fl(multiplier_accum, 1.0f);
136  float size_center = tempSize[0] * scalar;
137 
138  const int addXStepValue = QualityStepHelper::getStep();
139  const int addYStepValue = addXStepValue;
140  const int addXStepColor = addXStepValue * COM_DATA_TYPE_COLOR_CHANNELS;
141 
142  if (size_center > this->m_threshold) {
143  for (int ny = miny; ny < maxy; ny += addYStepValue) {
144  float dy = ny - y;
145  int offsetValueNy = ny * inputSizeBuffer->getWidth();
146  int offsetValueNxNy = offsetValueNy + (minx);
147  int offsetColorNxNy = offsetValueNxNy * COM_DATA_TYPE_COLOR_CHANNELS;
148  for (int nx = minx; nx < maxx; nx += addXStepValue) {
149  if (nx != x || ny != y) {
150  float size = MIN2(inputSizeFloatBuffer[offsetValueNxNy] * scalar, size_center);
151  if (size > this->m_threshold) {
152  float dx = nx - x;
153  if (size > fabsf(dx) && size > fabsf(dy)) {
154  float uv[2] = {
156  (dx / size) * (float)((COM_BLUR_BOKEH_PIXELS / 2) - 1),
158  (dy / size) * (float)((COM_BLUR_BOKEH_PIXELS / 2) - 1),
159  };
160  inputBokehBuffer->read(bokeh, uv[0], uv[1]);
161  madd_v4_v4v4(color_accum, bokeh, &inputProgramFloatBuffer[offsetColorNxNy]);
162  add_v4_v4(multiplier_accum, bokeh);
163  }
164  }
165  }
166  offsetColorNxNy += addXStepColor;
167  offsetValueNxNy += addXStepValue;
168  }
169  }
170  }
171 
172  output[0] = color_accum[0] / multiplier_accum[0];
173  output[1] = color_accum[1] / multiplier_accum[1];
174  output[2] = color_accum[2] / multiplier_accum[2];
175  output[3] = color_accum[3] / multiplier_accum[3];
176 
177  /* blend in out values over the threshold, otherwise we get sharp, ugly transitions */
178  if ((size_center > this->m_threshold) && (size_center < this->m_threshold * 2.0f)) {
179  /* factor from 0-1 */
180  float fac = (size_center - this->m_threshold) / this->m_threshold;
181  interp_v4_v4v4(output, readColor, output, fac);
182  }
183  }
184 }
185 
187  MemoryBuffer *outputMemoryBuffer,
188  cl_mem clOutputBuffer,
189  MemoryBuffer **inputMemoryBuffers,
190  std::list<cl_mem> *clMemToCleanUp,
191  std::list<cl_kernel> * /*clKernelsToCleanUp*/)
192 {
193  cl_kernel defocusKernel = device->COM_clCreateKernel("defocusKernel", nullptr);
194 
195  cl_int step = this->getStep();
196  cl_int maxBlur;
197  cl_float threshold = this->m_threshold;
198 
199  MemoryBuffer *sizeMemoryBuffer = this->m_inputSizeProgram->getInputMemoryBuffer(
200  inputMemoryBuffers);
201 
202  const float max_dim = MAX2(m_width, m_height);
203  cl_float scalar = this->m_do_size_scale ? (max_dim / 100.0f) : 1.0f;
204 
205  maxBlur = (cl_int)min_ff(sizeMemoryBuffer->get_max_value() * scalar, (float)this->m_maxBlur);
206 
208  defocusKernel, 0, -1, clMemToCleanUp, inputMemoryBuffers, this->m_inputProgram);
210  defocusKernel, 1, -1, clMemToCleanUp, inputMemoryBuffers, this->m_inputBokehProgram);
212  defocusKernel, 2, 4, clMemToCleanUp, inputMemoryBuffers, this->m_inputSizeProgram);
213  device->COM_clAttachOutputMemoryBufferToKernelParameter(defocusKernel, 3, clOutputBuffer);
214  device->COM_clAttachMemoryBufferOffsetToKernelParameter(defocusKernel, 5, outputMemoryBuffer);
215  clSetKernelArg(defocusKernel, 6, sizeof(cl_int), &step);
216  clSetKernelArg(defocusKernel, 7, sizeof(cl_int), &maxBlur);
217  clSetKernelArg(defocusKernel, 8, sizeof(cl_float), &threshold);
218  clSetKernelArg(defocusKernel, 9, sizeof(cl_float), &scalar);
219  device->COM_clAttachSizeToKernelParameter(defocusKernel, 10, this);
220 
221  device->COM_clEnqueueRange(defocusKernel, outputMemoryBuffer, 11, this);
222 }
223 
225 {
226  this->m_inputProgram = nullptr;
227  this->m_inputBokehProgram = nullptr;
228  this->m_inputSizeProgram = nullptr;
229 #ifdef COM_DEFOCUS_SEARCH
230  this->m_inputSearchProgram = nullptr;
231 #endif
232 }
233 
235  rcti *input, ReadBufferOperation *readOperation, rcti *output)
236 {
237  rcti newInput;
238  rcti bokehInput;
239 
240  const float max_dim = MAX2(m_width, m_height);
241  const float scalar = this->m_do_size_scale ? (max_dim / 100.0f) : 1.0f;
242  int maxBlurScalar = this->m_maxBlur * scalar;
243 
244  newInput.xmax = input->xmax + maxBlurScalar + 2;
245  newInput.xmin = input->xmin - maxBlurScalar + 2;
246  newInput.ymax = input->ymax + maxBlurScalar - 2;
247  newInput.ymin = input->ymin - maxBlurScalar - 2;
248  bokehInput.xmax = COM_BLUR_BOKEH_PIXELS;
249  bokehInput.xmin = 0;
250  bokehInput.ymax = COM_BLUR_BOKEH_PIXELS;
251  bokehInput.ymin = 0;
252 
253  NodeOperation *operation = getInputOperation(2);
254  if (operation->determineDependingAreaOfInterest(&newInput, readOperation, output)) {
255  return true;
256  }
257  operation = getInputOperation(1);
258  if (operation->determineDependingAreaOfInterest(&bokehInput, readOperation, output)) {
259  return true;
260  }
261 #ifdef COM_DEFOCUS_SEARCH
262  rcti searchInput;
263  searchInput.xmax = (input->xmax / InverseSearchRadiusOperation::DIVIDER) + 1;
264  searchInput.xmin = (input->xmin / InverseSearchRadiusOperation::DIVIDER) - 1;
265  searchInput.ymax = (input->ymax / InverseSearchRadiusOperation::DIVIDER) + 1;
266  searchInput.ymin = (input->ymin / InverseSearchRadiusOperation::DIVIDER) - 1;
267  operation = getInputOperation(3);
268  if (operation->determineDependingAreaOfInterest(&searchInput, readOperation, output)) {
269  return true;
270  }
271 #endif
272  operation = getInputOperation(0);
273  if (operation->determineDependingAreaOfInterest(&newInput, readOperation, output)) {
274  return true;
275  }
276  return false;
277 }
278 
279 #ifdef COM_DEFOCUS_SEARCH
280 // InverseSearchRadiusOperation
281 InverseSearchRadiusOperation::InverseSearchRadiusOperation()
282 {
283  this->addInputSocket(DataType::Value, ResizeMode::None); // radius
284  this->addOutputSocket(DataType::Color);
285  this->flags.complex = true;
286  this->m_inputRadius = nullptr;
287 }
288 
289 void InverseSearchRadiusOperation::initExecution()
290 {
291  this->m_inputRadius = this->getInputSocketReader(0);
292 }
293 
294 void *InverseSearchRadiusOperation::initializeTileData(rcti *rect)
295 {
296  MemoryBuffer *data = new MemoryBuffer(DataType::Color, rect);
297  float *buffer = data->getBuffer();
298  int x, y;
299  int width = this->m_inputRadius->getWidth();
300  int height = this->m_inputRadius->getHeight();
301  float temp[4];
302  int offset = 0;
303  for (y = rect->ymin; y < rect->ymax; y++) {
304  for (x = rect->xmin; x < rect->xmax; x++) {
305  int rx = x * DIVIDER;
306  int ry = y * DIVIDER;
307  buffer[offset] = MAX2(rx - m_maxBlur, 0);
308  buffer[offset + 1] = MAX2(ry - m_maxBlur, 0);
309  buffer[offset + 2] = MIN2(rx + DIVIDER + m_maxBlur, width);
310  buffer[offset + 3] = MIN2(ry + DIVIDER + m_maxBlur, height);
311  offset += 4;
312  }
313  }
314  // for (x = rect->xmin; x < rect->xmax ; x++) {
315  // for (y = rect->ymin; y < rect->ymax ; y++) {
316  // int rx = x * DIVIDER;
317  // int ry = y * DIVIDER;
318  // float radius = 0.0f;
319  // float maxx = x;
320  // float maxy = y;
321 
322  // for (int x2 = 0 ; x2 < DIVIDER ; x2 ++) {
323  // for (int y2 = 0 ; y2 < DIVIDER ; y2 ++) {
324  // this->m_inputRadius->read(temp, rx+x2, ry+y2, PixelSampler::Nearest);
325  // if (radius < temp[0]) {
326  // radius = temp[0];
327  // maxx = x2;
328  // maxy = y2;
329  // }
330  // }
331  // }
332  // int impactRadius = ceil(radius / DIVIDER);
333  // for (int x2 = x - impactRadius ; x2 < x + impactRadius ; x2 ++) {
334  // for (int y2 = y - impactRadius ; y2 < y + impactRadius ; y2 ++) {
335  // data->read(temp, x2, y2);
336  // temp[0] = MIN2(temp[0], maxx);
337  // temp[1] = MIN2(temp[1], maxy);
338  // temp[2] = MAX2(temp[2], maxx);
339  // temp[3] = MAX2(temp[3], maxy);
340  // data->writePixel(x2, y2, temp);
341  // }
342  // }
343  // }
344  // }
345  return data;
346 }
347 
348 void InverseSearchRadiusOperation::executePixelChunk(float output[4], int x, int y, void *data)
349 {
350  MemoryBuffer *buffer = (MemoryBuffer *)data;
351  buffer->readNoCheck(output, x, y);
352 }
353 
354 void InverseSearchRadiusOperation::deinitializeTileData(rcti *rect, void *data)
355 {
356  if (data) {
357  MemoryBuffer *mb = (MemoryBuffer *)data;
358  delete mb;
359  }
360 }
361 
362 void InverseSearchRadiusOperation::deinitExecution()
363 {
364  this->m_inputRadius = nullptr;
365 }
366 
367 void InverseSearchRadiusOperation::determineResolution(unsigned int resolution[2],
368  unsigned int preferredResolution[2])
369 {
370  NodeOperation::determineResolution(resolution, preferredResolution);
371  resolution[0] = resolution[0] / DIVIDER;
372  resolution[1] = resolution[1] / DIVIDER;
373 }
374 
375 bool InverseSearchRadiusOperation::determineDependingAreaOfInterest(
376  rcti *input, ReadBufferOperation *readOperation, rcti *output)
377 {
378  rcti newRect;
379  newRect.ymin = input->ymin * DIVIDER - m_maxBlur;
380  newRect.ymax = input->ymax * DIVIDER + m_maxBlur;
381  newRect.xmin = input->xmin * DIVIDER - m_maxBlur;
382  newRect.xmax = input->xmax * DIVIDER + m_maxBlur;
383  return NodeOperation::determineDependingAreaOfInterest(&newRect, readOperation, output);
384 }
385 #endif
386 
387 } // namespace blender::compositor
typedef float(TangentPoint)[2]
#define BLI_assert(a)
Definition: BLI_assert.h:58
MINLINE float min_ff(float a, float b)
MINLINE void add_v4_v4(float r[4], const float a[4])
MINLINE void copy_v4_v4(float r[4], const float a[4])
void interp_v4_v4v4(float r[4], const float a[4], const float b[4], const float t)
Definition: math_vector.c:58
MINLINE void madd_v4_v4v4(float r[4], const float a[4], const float b[4])
MINLINE void copy_v4_fl(float r[4], float f)
#define MAX2(a, b)
#define MIN2(a, b)
_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 const void *pointer _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_BOOL GLfloat param _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID GLenum GLfloat param _GL_VOID_RET _GL_VOID GLenum GLint param _GL_VOID_RET _GL_VOID GLushort pattern _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLint const GLdouble *points _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLint GLdouble GLdouble GLint GLint const GLdouble *points _GL_VOID_RET _GL_VOID GLdouble GLdouble u2 _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLdouble GLdouble v2 _GL_VOID_RET _GL_VOID GLenum GLfloat param _GL_VOID_RET _GL_VOID GLenum GLint param _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLdouble ny
_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 width
_GL_VOID GLfloat value _GL_VOID_RET _GL_VOID const GLuint GLboolean *residences _GL_BOOL_RET _GL_VOID GLsizei height
_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
Group RGB to Bright Vector Camera CLAMP
static DBVT_INLINE btScalar size(const btDbvtVolume &a)
Definition: btDbvt.cpp:52
#define output
a MemoryBuffer contains access to the data of a chunk
void read(float *result, int x, int y, MemoryBufferExtend extend_x=MemoryBufferExtend::Clip, MemoryBufferExtend extend_y=MemoryBufferExtend::Clip)
void readNoCheck(float *result, int x, int y, MemoryBufferExtend extend_x=MemoryBufferExtend::Clip, MemoryBufferExtend extend_y=MemoryBufferExtend::Clip)
const int getHeight() const
get the height of this MemoryBuffer
const int getWidth() const
get the width of this MemoryBuffer
float * getBuffer()
get the data of this MemoryBuffer
NodeOperation contains calculation logic.
virtual void * initializeTileData(rcti *)
void addInputSocket(DataType datatype, ResizeMode resize_mode=ResizeMode::Center)
NodeOperation * getInputOperation(unsigned int inputSocketindex)
virtual MemoryBuffer * getInputMemoryBuffer(MemoryBuffer **)
void addOutputSocket(DataType datatype)
virtual void determineResolution(unsigned int resolution[2], unsigned int preferredResolution[2])
determine the resolution of this node
SocketReader * getInputSocketReader(unsigned int inputSocketindex)
virtual bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output)
device representing an GPU OpenCL device. an instance of this class represents a single cl_device
void COM_clAttachMemoryBufferOffsetToKernelParameter(cl_kernel kernel, int offsetIndex, MemoryBuffer *memoryBuffers)
cl_mem COM_clAttachMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, int offsetIndex, std::list< cl_mem > *cleanup, MemoryBuffer **inputMemoryBuffers, SocketReader *reader)
void COM_clEnqueueRange(cl_kernel kernel, MemoryBuffer *outputMemoryBuffer)
void COM_clAttachOutputMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, cl_mem clOutputMemoryBuffer)
void COM_clAttachSizeToKernelParameter(cl_kernel kernel, int offsetIndex, NodeOperation *operation)
cl_kernel COM_clCreateKernel(const char *kernelname, std::list< cl_kernel > *clKernelsToCleanUp)
bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output) override
void executePixel(float output[4], int x, int y, void *data) override
void executeOpenCL(OpenCLDevice *device, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, std::list< cl_mem > *clMemToCleanUp, std::list< cl_kernel > *clKernelsToCleanUp) override
custom handle to add new tasks to the OpenCL command queue in order to execute a chunk on an GPUDevic...
@ None
The bottom left of the input image is the bottom left of the working area of the node,...
#define fabsf(x)
__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
constexpr int COM_DATA_TYPE_COLOR_CHANNELS
Definition: COM_defines.h:53
constexpr float COM_BLUR_BOKEH_PIXELS
Definition: COM_defines.h:80
int ymin
Definition: DNA_vec_types.h:80
int ymax
Definition: DNA_vec_types.h:80
int xmin
Definition: DNA_vec_types.h:79
int xmax
Definition: DNA_vec_types.h:79