Blender  V2.93
COM_WriteBufferOperation.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 "COM_OpenCLDevice.h"
21 #include "COM_defines.h"
22 #include <cstdio>
23 
24 namespace blender::compositor {
25 
27 {
28  this->addInputSocket(datatype);
29  this->m_memoryProxy = new MemoryProxy(datatype);
30  this->m_memoryProxy->setWriteBufferOperation(this);
31  this->m_memoryProxy->setExecutor(nullptr);
33 }
35 {
36  if (this->m_memoryProxy) {
37  delete this->m_memoryProxy;
38  this->m_memoryProxy = nullptr;
39  }
40 }
41 
43  float x,
44  float y,
45  PixelSampler sampler)
46 {
47  this->m_input->readSampled(output, x, y, sampler);
48 }
49 
51 {
52  this->m_input = this->getInputOperation(0);
53  this->m_memoryProxy->allocate(this->m_width, this->m_height);
54 }
55 
57 {
58  this->m_input = nullptr;
59  this->m_memoryProxy->free();
60 }
61 
62 void WriteBufferOperation::executeRegion(rcti *rect, unsigned int /*tileNumber*/)
63 {
64  MemoryBuffer *memoryBuffer = this->m_memoryProxy->getBuffer();
65  float *buffer = memoryBuffer->getBuffer();
66  const uint8_t num_channels = memoryBuffer->get_num_channels();
67  if (this->m_input->get_flags().complex) {
68  void *data = this->m_input->initializeTileData(rect);
69  int x1 = rect->xmin;
70  int y1 = rect->ymin;
71  int x2 = rect->xmax;
72  int y2 = rect->ymax;
73  int x;
74  int y;
75  bool breaked = false;
76  for (y = y1; y < y2 && (!breaked); y++) {
77  int offset4 = (y * memoryBuffer->getWidth() + x1) * num_channels;
78  for (x = x1; x < x2; x++) {
79  this->m_input->read(&(buffer[offset4]), x, y, data);
80  offset4 += num_channels;
81  }
82  if (isBraked()) {
83  breaked = true;
84  }
85  }
86  if (data) {
87  this->m_input->deinitializeTileData(rect, data);
88  data = nullptr;
89  }
90  }
91  else {
92  int x1 = rect->xmin;
93  int y1 = rect->ymin;
94  int x2 = rect->xmax;
95  int y2 = rect->ymax;
96 
97  int x;
98  int y;
99  bool breaked = false;
100  for (y = y1; y < y2 && (!breaked); y++) {
101  int offset4 = (y * memoryBuffer->getWidth() + x1) * num_channels;
102  for (x = x1; x < x2; x++) {
103  this->m_input->readSampled(&(buffer[offset4]), x, y, PixelSampler::Nearest);
104  offset4 += num_channels;
105  }
106  if (isBraked()) {
107  breaked = true;
108  }
109  }
110  }
111 }
112 
114  rcti * /*rect*/,
115  unsigned int /*chunkNumber*/,
116  MemoryBuffer **inputMemoryBuffers,
117  MemoryBuffer *outputBuffer)
118 {
119  float *outputFloatBuffer = outputBuffer->getBuffer();
120  cl_int error;
121  /*
122  * 1. create cl_mem from outputbuffer
123  * 2. call NodeOperation (input) executeOpenCLChunk(.....)
124  * 3. schedule read back from opencl to main device (outputbuffer)
125  * 4. schedule native callback
126  *
127  * note: list of cl_mem will be filled by 2, and needs to be cleaned up by 4
128  */
129  // STEP 1
130  const unsigned int outputBufferWidth = outputBuffer->getWidth();
131  const unsigned int outputBufferHeight = outputBuffer->getHeight();
132 
133  const cl_image_format *imageFormat = OpenCLDevice::determineImageFormat(outputBuffer);
134 
135  cl_mem clOutputBuffer = clCreateImage2D(device->getContext(),
136  CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
137  imageFormat,
138  outputBufferWidth,
139  outputBufferHeight,
140  0,
141  outputFloatBuffer,
142  &error);
143  if (error != CL_SUCCESS) {
144  printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
145  }
146 
147  // STEP 2
148  std::list<cl_mem> *clMemToCleanUp = new std::list<cl_mem>();
149  clMemToCleanUp->push_back(clOutputBuffer);
150  std::list<cl_kernel> *clKernelsToCleanUp = new std::list<cl_kernel>();
151 
152  this->m_input->executeOpenCL(device,
153  outputBuffer,
154  clOutputBuffer,
155  inputMemoryBuffers,
156  clMemToCleanUp,
157  clKernelsToCleanUp);
158 
159  // STEP 3
160 
161  size_t origin[3] = {0, 0, 0};
162  size_t region[3] = {outputBufferWidth, outputBufferHeight, 1};
163 
164  // clFlush(queue);
165  // clFinish(queue);
166 
167  error = clEnqueueBarrier(device->getQueue());
168  if (error != CL_SUCCESS) {
169  printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
170  }
171  error = clEnqueueReadImage(device->getQueue(),
172  clOutputBuffer,
173  CL_TRUE,
174  origin,
175  region,
176  0,
177  0,
178  outputFloatBuffer,
179  0,
180  nullptr,
181  nullptr);
182  if (error != CL_SUCCESS) {
183  printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
184  }
185 
186  this->getMemoryProxy()->getBuffer()->fill_from(*outputBuffer);
187 
188  // STEP 4
189  while (!clMemToCleanUp->empty()) {
190  cl_mem mem = clMemToCleanUp->front();
191  error = clReleaseMemObject(mem);
192  if (error != CL_SUCCESS) {
193  printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
194  }
195  clMemToCleanUp->pop_front();
196  }
197 
198  while (!clKernelsToCleanUp->empty()) {
199  cl_kernel kernel = clKernelsToCleanUp->front();
200  error = clReleaseKernel(kernel);
201  if (error != CL_SUCCESS) {
202  printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
203  }
204  clKernelsToCleanUp->pop_front();
205  }
206  delete clKernelsToCleanUp;
207 }
208 
209 void WriteBufferOperation::determineResolution(unsigned int resolution[2],
210  unsigned int preferredResolution[2])
211 {
212  NodeOperation::determineResolution(resolution, preferredResolution);
213  /* make sure there is at least one pixel stored in case the input is a single value */
214  m_single_value = false;
215  if (resolution[0] == 0) {
216  resolution[0] = 1;
217  m_single_value = true;
218  }
219  if (resolution[1] == 0) {
220  resolution[1] = 1;
221  m_single_value = true;
222  }
223 }
224 
226 {
227  NodeOperation *inputOperation = this->getInputOperation(0);
228  this->setWidth(inputOperation->getWidth());
229  this->setHeight(inputOperation->getHeight());
230 }
231 
232 } // namespace blender::compositor
_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 GLdouble nz _GL_VOID_RET _GL_VOID GLfloat GLfloat nz _GL_VOID_RET _GL_VOID GLint GLint nz _GL_VOID_RET _GL_VOID GLshort GLshort nz _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_VOID GLsizei const GLfloat *values _GL_VOID_RET _GL_VOID GLsizei const GLushort *values _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID const GLuint const GLclampf *priorities _GL_VOID_RET _GL_VOID GLdouble y _GL_VOID_RET _GL_VOID GLfloat y _GL_VOID_RET _GL_VOID GLint y _GL_VOID_RET _GL_VOID GLshort y _GL_VOID_RET _GL_VOID GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLfloat GLfloat z _GL_VOID_RET _GL_VOID GLint GLint z _GL_VOID_RET _GL_VOID GLshort GLshort z _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble w _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat w _GL_VOID_RET _GL_VOID GLint GLint GLint w _GL_VOID_RET _GL_VOID GLshort GLshort GLshort w _GL_VOID_RET _GL_VOID GLdouble y1
_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 GLdouble nz _GL_VOID_RET _GL_VOID GLfloat GLfloat nz _GL_VOID_RET _GL_VOID GLint GLint nz _GL_VOID_RET _GL_VOID GLshort GLshort nz _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_VOID GLsizei const GLfloat *values _GL_VOID_RET _GL_VOID GLsizei const GLushort *values _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID const GLuint const GLclampf *priorities _GL_VOID_RET _GL_VOID GLdouble y _GL_VOID_RET _GL_VOID GLfloat y _GL_VOID_RET _GL_VOID GLint y _GL_VOID_RET _GL_VOID GLshort y _GL_VOID_RET _GL_VOID GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLfloat GLfloat z _GL_VOID_RET _GL_VOID GLint GLint z _GL_VOID_RET _GL_VOID GLshort GLshort z _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble w _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat w _GL_VOID_RET _GL_VOID GLint GLint GLint w _GL_VOID_RET _GL_VOID GLshort GLshort GLshort w _GL_VOID_RET _GL_VOID GLdouble GLdouble x2
_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
#define output
a MemoryBuffer contains access to the data of a chunk
const int getHeight() const
get the height of this MemoryBuffer
const int getWidth() const
get the width of this MemoryBuffer
void fill_from(const MemoryBuffer &src)
add the content from otherBuffer to this MemoryBuffer
float * getBuffer()
get the data of this MemoryBuffer
A MemoryProxy is a unique identifier for a memory buffer. A single MemoryProxy is used among all chun...
void free()
free the allocated memory
void setExecutor(ExecutionGroup *executor)
set the ExecutionGroup that can be scheduled to calculate a certain chunk.
MemoryBuffer * getBuffer()
get the allocated memory
void allocate(unsigned int width, unsigned int height)
allocate memory of size width x height
void setWriteBufferOperation(WriteBufferOperation *operation)
set the WriteBufferOperation that is responsible for writing to this MemoryProxy
NodeOperation contains calculation logic.
virtual void * initializeTileData(rcti *)
void readSampled(float result[4], float x, float y, PixelSampler sampler)
void addInputSocket(DataType datatype, ResizeMode resize_mode=ResizeMode::Center)
NodeOperation * getInputOperation(unsigned int inputSocketindex)
void read(float result[4], int x, int y, void *chunkData)
virtual void executeOpenCL(OpenCLDevice *, MemoryBuffer *, cl_mem, MemoryBuffer **, std::list< cl_mem > *, std::list< cl_kernel > *)
custom handle to add new tasks to the OpenCL command queue in order to execute a chunk on an GPUDevic...
const NodeOperationFlags get_flags() const
virtual void deinitializeTileData(rcti *, void *)
virtual void determineResolution(unsigned int resolution[2], unsigned int preferredResolution[2])
determine the resolution of this node
void setWidth(unsigned int width)
void setHeight(unsigned int height)
device representing an GPU OpenCL device. an instance of this class represents a single cl_device
static const cl_image_format * determineImageFormat(MemoryBuffer *memoryBuffer)
determine an image format
void executeRegion(rcti *rect, unsigned int tileNumber) override
when a chunk is executed by a CPUDevice, this method is called
void executeOpenCLRegion(OpenCLDevice *device, rcti *rect, unsigned int chunkNumber, MemoryBuffer **memoryBuffers, MemoryBuffer *outputBuffer) override
when a chunk is executed by an OpenCLDevice, this method is called
void executePixelSampled(float output[4], float x, float y, PixelSampler sampler) override
calculate a single pixel
void determineResolution(unsigned int resolution[2], unsigned int preferredResolution[2]) override
determine the resolution of this node
DataType
possible data types for sockets
Definition: COM_defines.h:27
__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
static void error(const char *str)
Definition: meshlaplacian.c:65
unsigned char uint8_t
Definition: stdint.h:81
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