Blender  V2.93
COM_DilateErodeOperation.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 "MEM_guardedalloc.h"
24 
25 namespace blender::compositor {
26 
27 // DilateErode Distance Threshold
29 {
32  this->flags.complex = true;
33  this->m_inputProgram = nullptr;
34  this->m_inset = 0.0f;
35  this->m__switch = 0.5f;
36  this->m_distance = 0.0f;
37 }
39 {
40  this->m_inputProgram = this->getInputSocketReader(0);
41  if (this->m_distance < 0.0f) {
42  this->m_scope = -this->m_distance + this->m_inset;
43  }
44  else {
45  if (this->m_inset * 2 > this->m_distance) {
46  this->m_scope = MAX2(this->m_inset * 2 - this->m_distance, this->m_distance);
47  }
48  else {
49  this->m_scope = this->m_distance;
50  }
51  }
52  if (this->m_scope < 3) {
53  this->m_scope = 3;
54  }
55 }
56 
58 {
59  void *buffer = this->m_inputProgram->initializeTileData(nullptr);
60  return buffer;
61 }
62 
64 {
65  float inputValue[4];
66  const float sw = this->m__switch;
67  const float distance = this->m_distance;
68  float pixelvalue;
69  const float rd = this->m_scope * this->m_scope;
70  const float inset = this->m_inset;
71  float mindist = rd * 2;
72 
73  MemoryBuffer *inputBuffer = (MemoryBuffer *)data;
74  float *buffer = inputBuffer->getBuffer();
75  const rcti &input_rect = inputBuffer->get_rect();
76  const int minx = MAX2(x - this->m_scope, input_rect.xmin);
77  const int miny = MAX2(y - this->m_scope, input_rect.ymin);
78  const int maxx = MIN2(x + this->m_scope, input_rect.xmax);
79  const int maxy = MIN2(y + this->m_scope, input_rect.ymax);
80  const int bufferWidth = inputBuffer->getWidth();
81  int offset;
82 
83  inputBuffer->read(inputValue, x, y);
84  if (inputValue[0] > sw) {
85  for (int yi = miny; yi < maxy; yi++) {
86  const float dy = yi - y;
87  offset = ((yi - input_rect.ymin) * bufferWidth + (minx - input_rect.xmin));
88  for (int xi = minx; xi < maxx; xi++) {
89  if (buffer[offset] < sw) {
90  const float dx = xi - x;
91  const float dis = dx * dx + dy * dy;
92  mindist = MIN2(mindist, dis);
93  }
94  offset++;
95  }
96  }
97  pixelvalue = -sqrtf(mindist);
98  }
99  else {
100  for (int yi = miny; yi < maxy; yi++) {
101  const float dy = yi - y;
102  offset = ((yi - input_rect.ymin) * bufferWidth + (minx - input_rect.xmin));
103  for (int xi = minx; xi < maxx; xi++) {
104  if (buffer[offset] > sw) {
105  const float dx = xi - x;
106  const float dis = dx * dx + dy * dy;
107  mindist = MIN2(mindist, dis);
108  }
109  offset++;
110  }
111  }
112  pixelvalue = sqrtf(mindist);
113  }
114 
115  if (distance > 0.0f) {
116  const float delta = distance - pixelvalue;
117  if (delta >= 0.0f) {
118  if (delta >= inset) {
119  output[0] = 1.0f;
120  }
121  else {
122  output[0] = delta / inset;
123  }
124  }
125  else {
126  output[0] = 0.0f;
127  }
128  }
129  else {
130  const float delta = -distance + pixelvalue;
131  if (delta < 0.0f) {
132  if (delta < -inset) {
133  output[0] = 1.0f;
134  }
135  else {
136  output[0] = (-delta) / inset;
137  }
138  }
139  else {
140  output[0] = 0.0f;
141  }
142  }
143 }
144 
146 {
147  this->m_inputProgram = nullptr;
148 }
149 
151  rcti *input, ReadBufferOperation *readOperation, rcti *output)
152 {
153  rcti newInput;
154 
155  newInput.xmax = input->xmax + this->m_scope;
156  newInput.xmin = input->xmin - this->m_scope;
157  newInput.ymax = input->ymax + this->m_scope;
158  newInput.ymin = input->ymin - this->m_scope;
159 
160  return NodeOperation::determineDependingAreaOfInterest(&newInput, readOperation, output);
161 }
162 
163 // Dilate Distance
165 {
168  this->m_inputProgram = nullptr;
169  this->m_distance = 0.0f;
170  flags.complex = true;
171  flags.open_cl = true;
172 }
174 {
175  this->m_inputProgram = this->getInputSocketReader(0);
176  this->m_scope = this->m_distance;
177  if (this->m_scope < 3) {
178  this->m_scope = 3;
179  }
180 }
181 
183 {
184  void *buffer = this->m_inputProgram->initializeTileData(nullptr);
185  return buffer;
186 }
187 
188 void DilateDistanceOperation::executePixel(float output[4], int x, int y, void *data)
189 {
190  const float distance = this->m_distance;
191  const float mindist = distance * distance;
192 
193  MemoryBuffer *inputBuffer = (MemoryBuffer *)data;
194  float *buffer = inputBuffer->getBuffer();
195  const rcti &input_rect = inputBuffer->get_rect();
196  const int minx = MAX2(x - this->m_scope, input_rect.xmin);
197  const int miny = MAX2(y - this->m_scope, input_rect.ymin);
198  const int maxx = MIN2(x + this->m_scope, input_rect.xmax);
199  const int maxy = MIN2(y + this->m_scope, input_rect.ymax);
200  const int bufferWidth = inputBuffer->getWidth();
201  int offset;
202 
203  float value = 0.0f;
204 
205  for (int yi = miny; yi < maxy; yi++) {
206  const float dy = yi - y;
207  offset = ((yi - input_rect.ymin) * bufferWidth + (minx - input_rect.xmin));
208  for (int xi = minx; xi < maxx; xi++) {
209  const float dx = xi - x;
210  const float dis = dx * dx + dy * dy;
211  if (dis <= mindist) {
212  value = MAX2(buffer[offset], value);
213  }
214  offset++;
215  }
216  }
217  output[0] = value;
218 }
219 
221 {
222  this->m_inputProgram = nullptr;
223 }
224 
226  ReadBufferOperation *readOperation,
227  rcti *output)
228 {
229  rcti newInput;
230 
231  newInput.xmax = input->xmax + this->m_scope;
232  newInput.xmin = input->xmin - this->m_scope;
233  newInput.ymax = input->ymax + this->m_scope;
234  newInput.ymin = input->ymin - this->m_scope;
235 
236  return NodeOperation::determineDependingAreaOfInterest(&newInput, readOperation, output);
237 }
238 
240  MemoryBuffer *outputMemoryBuffer,
241  cl_mem clOutputBuffer,
242  MemoryBuffer **inputMemoryBuffers,
243  std::list<cl_mem> *clMemToCleanUp,
244  std::list<cl_kernel> * /*clKernelsToCleanUp*/)
245 {
246  cl_kernel dilateKernel = device->COM_clCreateKernel("dilateKernel", nullptr);
247 
248  cl_int distanceSquared = this->m_distance * this->m_distance;
249  cl_int scope = this->m_scope;
250 
252  dilateKernel, 0, 2, clMemToCleanUp, inputMemoryBuffers, this->m_inputProgram);
253  device->COM_clAttachOutputMemoryBufferToKernelParameter(dilateKernel, 1, clOutputBuffer);
254  device->COM_clAttachMemoryBufferOffsetToKernelParameter(dilateKernel, 3, outputMemoryBuffer);
255  clSetKernelArg(dilateKernel, 4, sizeof(cl_int), &scope);
256  clSetKernelArg(dilateKernel, 5, sizeof(cl_int), &distanceSquared);
257  device->COM_clAttachSizeToKernelParameter(dilateKernel, 6, this);
258  device->COM_clEnqueueRange(dilateKernel, outputMemoryBuffer, 7, this);
259 }
260 
261 // Erode Distance
263 {
264  /* pass */
265 }
266 
267 void ErodeDistanceOperation::executePixel(float output[4], int x, int y, void *data)
268 {
269  const float distance = this->m_distance;
270  const float mindist = distance * distance;
271 
272  MemoryBuffer *inputBuffer = (MemoryBuffer *)data;
273  float *buffer = inputBuffer->getBuffer();
274  const rcti &input_rect = inputBuffer->get_rect();
275  const int minx = MAX2(x - this->m_scope, input_rect.xmin);
276  const int miny = MAX2(y - this->m_scope, input_rect.ymin);
277  const int maxx = MIN2(x + this->m_scope, input_rect.xmax);
278  const int maxy = MIN2(y + this->m_scope, input_rect.ymax);
279  const int bufferWidth = inputBuffer->getWidth();
280  int offset;
281 
282  float value = 1.0f;
283 
284  for (int yi = miny; yi < maxy; yi++) {
285  const float dy = yi - y;
286  offset = ((yi - input_rect.ymin) * bufferWidth + (minx - input_rect.xmin));
287  for (int xi = minx; xi < maxx; xi++) {
288  const float dx = xi - x;
289  const float dis = dx * dx + dy * dy;
290  if (dis <= mindist) {
291  value = MIN2(buffer[offset], value);
292  }
293  offset++;
294  }
295  }
296  output[0] = value;
297 }
298 
300  MemoryBuffer *outputMemoryBuffer,
301  cl_mem clOutputBuffer,
302  MemoryBuffer **inputMemoryBuffers,
303  std::list<cl_mem> *clMemToCleanUp,
304  std::list<cl_kernel> * /*clKernelsToCleanUp*/)
305 {
306  cl_kernel erodeKernel = device->COM_clCreateKernel("erodeKernel", nullptr);
307 
308  cl_int distanceSquared = this->m_distance * this->m_distance;
309  cl_int scope = this->m_scope;
310 
312  erodeKernel, 0, 2, clMemToCleanUp, inputMemoryBuffers, this->m_inputProgram);
313  device->COM_clAttachOutputMemoryBufferToKernelParameter(erodeKernel, 1, clOutputBuffer);
314  device->COM_clAttachMemoryBufferOffsetToKernelParameter(erodeKernel, 3, outputMemoryBuffer);
315  clSetKernelArg(erodeKernel, 4, sizeof(cl_int), &scope);
316  clSetKernelArg(erodeKernel, 5, sizeof(cl_int), &distanceSquared);
317  device->COM_clAttachSizeToKernelParameter(erodeKernel, 6, this);
318  device->COM_clEnqueueRange(erodeKernel, outputMemoryBuffer, 7, this);
319 }
320 
321 // Dilate step
323 {
326  this->flags.complex = true;
327  this->m_inputProgram = nullptr;
328 }
330 {
331  this->m_inputProgram = this->getInputSocketReader(0);
332 }
333 
334 // small helper to pass data from initializeTileData to executePixel
335 struct tile_info {
337  int width;
338  float *buffer;
339 };
340 
341 static tile_info *create_cache(int xmin, int xmax, int ymin, int ymax)
342 {
343  tile_info *result = (tile_info *)MEM_mallocN(sizeof(tile_info), "dilate erode tile");
344  result->rect.xmin = xmin;
345  result->rect.xmax = xmax;
346  result->rect.ymin = ymin;
347  result->rect.ymax = ymax;
348  result->width = xmax - xmin;
349  result->buffer = (float *)MEM_callocN(sizeof(float) * (ymax - ymin) * result->width,
350  "dilate erode cache");
351  return result;
352 }
353 
355 {
356  MemoryBuffer *tile = (MemoryBuffer *)this->m_inputProgram->initializeTileData(nullptr);
357  int x, y, i;
358  int width = tile->getWidth();
359  int height = tile->getHeight();
360  float *buffer = tile->getBuffer();
361 
362  int half_window = this->m_iterations;
363  int window = half_window * 2 + 1;
364 
365  int xmin = MAX2(0, rect->xmin - half_window);
366  int ymin = MAX2(0, rect->ymin - half_window);
367  int xmax = MIN2(width, rect->xmax + half_window);
368  int ymax = MIN2(height, rect->ymax + half_window);
369 
370  int bwidth = rect->xmax - rect->xmin;
371  int bheight = rect->ymax - rect->ymin;
372 
373  // Note: Cache buffer has original tilesize width, but new height.
374  // We have to calculate the additional rows in the first pass,
375  // to have valid data available for the second pass.
376  tile_info *result = create_cache(rect->xmin, rect->xmax, ymin, ymax);
377  float *rectf = result->buffer;
378 
379  // temp holds maxima for every step in the algorithm, buf holds a
380  // single row or column of input values, padded with FLT_MAX's to
381  // simplify the logic.
382  float *temp = (float *)MEM_mallocN(sizeof(float) * (2 * window - 1), "dilate erode temp");
383  float *buf = (float *)MEM_mallocN(sizeof(float) * (MAX2(bwidth, bheight) + 5 * half_window),
384  "dilate erode buf");
385 
386  // The following is based on the van Herk/Gil-Werman algorithm for morphology operations.
387  // first pass, horizontal dilate/erode
388  for (y = ymin; y < ymax; y++) {
389  for (x = 0; x < bwidth + 5 * half_window; x++) {
390  buf[x] = -FLT_MAX;
391  }
392  for (x = xmin; x < xmax; x++) {
393  buf[x - rect->xmin + window - 1] = buffer[(y * width + x)];
394  }
395 
396  for (i = 0; i < (bwidth + 3 * half_window) / window; i++) {
397  int start = (i + 1) * window - 1;
398 
399  temp[window - 1] = buf[start];
400  for (x = 1; x < window; x++) {
401  temp[window - 1 - x] = MAX2(temp[window - x], buf[start - x]);
402  temp[window - 1 + x] = MAX2(temp[window + x - 2], buf[start + x]);
403  }
404 
405  start = half_window + (i - 1) * window + 1;
406  for (x = -MIN2(0, start); x < window - MAX2(0, start + window - bwidth); x++) {
407  rectf[bwidth * (y - ymin) + (start + x)] = MAX2(temp[x], temp[x + window - 1]);
408  }
409  }
410  }
411 
412  // second pass, vertical dilate/erode
413  for (x = 0; x < bwidth; x++) {
414  for (y = 0; y < bheight + 5 * half_window; y++) {
415  buf[y] = -FLT_MAX;
416  }
417  for (y = ymin; y < ymax; y++) {
418  buf[y - rect->ymin + window - 1] = rectf[(y - ymin) * bwidth + x];
419  }
420 
421  for (i = 0; i < (bheight + 3 * half_window) / window; i++) {
422  int start = (i + 1) * window - 1;
423 
424  temp[window - 1] = buf[start];
425  for (y = 1; y < window; y++) {
426  temp[window - 1 - y] = MAX2(temp[window - y], buf[start - y]);
427  temp[window - 1 + y] = MAX2(temp[window + y - 2], buf[start + y]);
428  }
429 
430  start = half_window + (i - 1) * window + 1;
431  for (y = -MIN2(0, start); y < window - MAX2(0, start + window - bheight); y++) {
432  rectf[bwidth * (y + start + (rect->ymin - ymin)) + x] = MAX2(temp[y],
433  temp[y + window - 1]);
434  }
435  }
436  }
437 
438  MEM_freeN(temp);
439  MEM_freeN(buf);
440 
441  return result;
442 }
443 
444 void DilateStepOperation::executePixel(float output[4], int x, int y, void *data)
445 {
446  tile_info *tile = (tile_info *)data;
447  int nx = x - tile->rect.xmin;
448  int ny = y - tile->rect.ymin;
449  output[0] = tile->buffer[tile->width * ny + nx];
450 }
451 
453 {
454  this->m_inputProgram = nullptr;
455 }
456 
458 {
459  tile_info *tile = (tile_info *)data;
460  MEM_freeN(tile->buffer);
461  MEM_freeN(tile);
462 }
463 
465  ReadBufferOperation *readOperation,
466  rcti *output)
467 {
468  rcti newInput;
469  int it = this->m_iterations;
470  newInput.xmax = input->xmax + it;
471  newInput.xmin = input->xmin - it;
472  newInput.ymax = input->ymax + it;
473  newInput.ymin = input->ymin - it;
474 
475  return NodeOperation::determineDependingAreaOfInterest(&newInput, readOperation, output);
476 }
477 
478 // Erode step
480 {
481  /* pass */
482 }
483 
485 {
486  MemoryBuffer *tile = (MemoryBuffer *)this->m_inputProgram->initializeTileData(nullptr);
487  int x, y, i;
488  int width = tile->getWidth();
489  int height = tile->getHeight();
490  float *buffer = tile->getBuffer();
491 
492  int half_window = this->m_iterations;
493  int window = half_window * 2 + 1;
494 
495  int xmin = MAX2(0, rect->xmin - half_window);
496  int ymin = MAX2(0, rect->ymin - half_window);
497  int xmax = MIN2(width, rect->xmax + half_window);
498  int ymax = MIN2(height, rect->ymax + half_window);
499 
500  int bwidth = rect->xmax - rect->xmin;
501  int bheight = rect->ymax - rect->ymin;
502 
503  // Note: Cache buffer has original tilesize width, but new height.
504  // We have to calculate the additional rows in the first pass,
505  // to have valid data available for the second pass.
506  tile_info *result = create_cache(rect->xmin, rect->xmax, ymin, ymax);
507  float *rectf = result->buffer;
508 
509  // temp holds maxima for every step in the algorithm, buf holds a
510  // single row or column of input values, padded with FLT_MAX's to
511  // simplify the logic.
512  float *temp = (float *)MEM_mallocN(sizeof(float) * (2 * window - 1), "dilate erode temp");
513  float *buf = (float *)MEM_mallocN(sizeof(float) * (MAX2(bwidth, bheight) + 5 * half_window),
514  "dilate erode buf");
515 
516  // The following is based on the van Herk/Gil-Werman algorithm for morphology operations.
517  // first pass, horizontal dilate/erode
518  for (y = ymin; y < ymax; y++) {
519  for (x = 0; x < bwidth + 5 * half_window; x++) {
520  buf[x] = FLT_MAX;
521  }
522  for (x = xmin; x < xmax; x++) {
523  buf[x - rect->xmin + window - 1] = buffer[(y * width + x)];
524  }
525 
526  for (i = 0; i < (bwidth + 3 * half_window) / window; i++) {
527  int start = (i + 1) * window - 1;
528 
529  temp[window - 1] = buf[start];
530  for (x = 1; x < window; x++) {
531  temp[window - 1 - x] = MIN2(temp[window - x], buf[start - x]);
532  temp[window - 1 + x] = MIN2(temp[window + x - 2], buf[start + x]);
533  }
534 
535  start = half_window + (i - 1) * window + 1;
536  for (x = -MIN2(0, start); x < window - MAX2(0, start + window - bwidth); x++) {
537  rectf[bwidth * (y - ymin) + (start + x)] = MIN2(temp[x], temp[x + window - 1]);
538  }
539  }
540  }
541 
542  // second pass, vertical dilate/erode
543  for (x = 0; x < bwidth; x++) {
544  for (y = 0; y < bheight + 5 * half_window; y++) {
545  buf[y] = FLT_MAX;
546  }
547  for (y = ymin; y < ymax; y++) {
548  buf[y - rect->ymin + window - 1] = rectf[(y - ymin) * bwidth + x];
549  }
550 
551  for (i = 0; i < (bheight + 3 * half_window) / window; i++) {
552  int start = (i + 1) * window - 1;
553 
554  temp[window - 1] = buf[start];
555  for (y = 1; y < window; y++) {
556  temp[window - 1 - y] = MIN2(temp[window - y], buf[start - y]);
557  temp[window - 1 + y] = MIN2(temp[window + y - 2], buf[start + y]);
558  }
559 
560  start = half_window + (i - 1) * window + 1;
561  for (y = -MIN2(0, start); y < window - MAX2(0, start + window - bheight); y++) {
562  rectf[bwidth * (y + start + (rect->ymin - ymin)) + x] = MIN2(temp[y],
563  temp[y + window - 1]);
564  }
565  }
566  }
567 
568  MEM_freeN(temp);
569  MEM_freeN(buf);
570 
571  return result;
572 }
573 
574 } // namespace blender::compositor
#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
Read Guarded memory(de)allocation.
#define output
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...
void executePixel(float output[4], int x, int y, void *data) override
bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output) override
bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output) override
void executePixel(float output[4], int x, int y, void *data) override
bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output) override
void deinitializeTileData(rcti *rect, void *data) 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...
void executePixel(float output[4], int x, int y, void *data) override
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)
const rcti & get_rect() const
get the rect of this MemoryBuffer
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
virtual void * initializeTileData(rcti *)
void addInputSocket(DataType datatype, ResizeMode resize_mode=ResizeMode::Center)
void addOutputSocket(DataType datatype)
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)
#define sqrtf(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
void(* MEM_freeN)(void *vmemh)
Definition: mallocn.c:41
void *(* MEM_callocN)(size_t len, const char *str)
Definition: mallocn.c:45
void *(* MEM_mallocN)(size_t len, const char *str)
Definition: mallocn.c:47
static tile_info * create_cache(int xmin, int xmax, int ymin, int ymax)
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
ccl_device_inline float distance(const float2 &a, const float2 &b)