Medical Imaging Interaction Toolkit  2018.4.99-3e3f1a6e
Medical Imaging Interaction Toolkit
mitkOclFilter.cpp
Go to the documentation of this file.
1 /*============================================================================
2 
3 The Medical Imaging Interaction Toolkit (MITK)
4 
5 Copyright (c) German Cancer Research Center (DKFZ)
6 All rights reserved.
7 
8 Use of this source code is governed by a 3-clause BSD license that can be
9 found in the LICENSE file.
10 
11 ============================================================================*/
12 
13 //Ocl
14 #include "mitkOclFilter.h"
15 #include "mitkOclUtils.h"
16 #include "mitkOpenCLActivator.h"
17 
18 
19 //Mitk
20 #include <mitkLogMacros.h>
21 #include <mitkConfig.h>
22 
23 //usService
24 #include "usServiceReference.h"
25 #include <usServiceRegistration.h>
26 #include <usModuleContext.h>
27 #include <usGetModuleContext.h>
28 #include <usModule.h>
29 #include <usModuleResource.h>
30 #include <usModuleResourceStream.h>
31 
32 //standard library
33 #include <thread>
34 
36  : m_ClCompilerFlags(""),
37  m_ClProgram(nullptr),
38  m_CommandQue(nullptr),
39  m_FilterID("mitkOclFilter"),
40  m_Preambel(" "),
41  m_Initialized(false)
42 {
43 }
44 
45 mitk::OclFilter::OclFilter(const char* filename)
46  : m_ClCompilerFlags(""),
47  m_ClProgram(nullptr),
48  m_CommandQue(nullptr),
49  m_FilterID(filename),
50  m_Preambel(" "),
51  m_Initialized(false)
52 {
53  m_ClFiles.push_back(filename);
54 }
55 
57 {
58  MITK_DEBUG << "OclFilter Destructor";
59 
60  // release program
61  if (m_ClProgram)
62  {
65 
66  // remove program from storage
67  resources->RemoveProgram(m_FilterID);
68  }
69 }
70 
71 bool mitk::OclFilter::ExecuteKernel( cl_kernel kernel, unsigned int workSizeDim )
72 {
73  cl_int clErr = 0;
74 
75  clErr = clEnqueueNDRangeKernel( this->m_CommandQue, kernel, workSizeDim,
76  nullptr, this->m_GlobalWorkSize, m_LocalWorkSize, 0, nullptr, nullptr);
77 
78  CHECK_OCL_ERR( clErr );
79 
80  return ( clErr == CL_SUCCESS );
81 }
82 
83 bool mitk::OclFilter::ExecuteKernelChunks( cl_kernel kernel, unsigned int workSizeDim, size_t* chunksDim )
84 {
85  size_t offset[3] ={0, 0, 0};
86  cl_int clErr = 0;
87 
88  if(workSizeDim == 2)
89  {
90  for(offset[0] = 0; offset[0] < m_GlobalWorkSize[0]; offset[0] += chunksDim[0])
91  {
92  for(offset[1] = 0; offset[1] < m_GlobalWorkSize[1]; offset[1] += chunksDim[1])
93  {
94  clErr |= clEnqueueNDRangeKernel(this->m_CommandQue, kernel, workSizeDim,
95  offset, chunksDim, m_LocalWorkSize, 0, nullptr, nullptr);
96  }
97  }
98  }
99  else if(workSizeDim == 3)
100  {
101  for(offset[0] = 0; offset[0] < m_GlobalWorkSize[0]; offset[0] += chunksDim[0])
102  {
103  for(offset[1] = 0; offset[1] < m_GlobalWorkSize[1]; offset[1] += chunksDim[1])
104  {
105  for(offset[2] = 0; offset[2] < m_GlobalWorkSize[2]; offset[2] += chunksDim[2])
106  {
107  clErr |= clEnqueueNDRangeKernel( this->m_CommandQue, kernel, workSizeDim,
108  offset, chunksDim, m_LocalWorkSize, 0, nullptr, nullptr);
109  }
110  }
111  }
112  }
113 
114  CHECK_OCL_ERR(clErr);
115 
116  return ( clErr == CL_SUCCESS );
117 }
118 
119 bool mitk::OclFilter::ExecuteKernelChunksInBatches(cl_kernel kernel, unsigned int workSizeDim, size_t* chunksDim, size_t batchSize, int waitTimems)
120 {
121  size_t offset[3] = { 0, 0, 0 };
122  cl_int clErr = 0;
123 
124  unsigned int currentChunk = 0;
125  cl_event* waitFor = new cl_event[batchSize];
126 
127  if (workSizeDim == 2)
128  {
129  for (offset[0] = 0; offset[0] < m_GlobalWorkSize[0]; offset[0] += chunksDim[0])
130  {
131  for (offset[1] = 0; offset[1] < m_GlobalWorkSize[1]; offset[1] += chunksDim[1])
132  {
133  if (currentChunk % batchSize == 0 && currentChunk != 0)
134  {
135  clWaitForEvents(batchSize, &waitFor[0]);
136  std::this_thread::sleep_for(std::chrono::milliseconds(waitTimems));
137  clErr |= clEnqueueNDRangeKernel(this->m_CommandQue, kernel, workSizeDim,
138  offset, chunksDim, m_LocalWorkSize, 0, nullptr, &waitFor[0]);
139  }
140  else
141  {
142  clErr |= clEnqueueNDRangeKernel(this->m_CommandQue, kernel, workSizeDim,
143  offset, chunksDim, m_LocalWorkSize, 0, nullptr, &waitFor[currentChunk % batchSize]);
144  }
145  currentChunk++;
146  }
147  }
148  }
149  else if (workSizeDim == 3)
150  {
151  for (offset[0] = 0; offset[0] < m_GlobalWorkSize[0]; offset[0] += chunksDim[0])
152  {
153  for (offset[1] = 0; offset[1] < m_GlobalWorkSize[1]; offset[1] += chunksDim[1])
154  {
155  for (offset[2] = 0; offset[2] < m_GlobalWorkSize[2]; offset[2] += chunksDim[2])
156  {
157  if (currentChunk % batchSize == 0 && currentChunk != 0)
158  {
159  clWaitForEvents(batchSize, &waitFor[0]);
160  std::this_thread::sleep_for(std::chrono::milliseconds(waitTimems));
161  clErr |= clEnqueueNDRangeKernel(this->m_CommandQue, kernel, workSizeDim,
162  offset, chunksDim, m_LocalWorkSize, 0, nullptr, &waitFor[0]);
163  }
164  else
165  {
166  clErr |= clEnqueueNDRangeKernel(this->m_CommandQue, kernel, workSizeDim,
167  offset, chunksDim, m_LocalWorkSize, 0, nullptr, &waitFor[currentChunk % batchSize]);
168  }
169  currentChunk++;
170  }
171  }
172  }
173  }
174  CHECK_OCL_ERR(clErr);
175 
176  return (clErr == CL_SUCCESS);
177 }
178 
179 
181 {
184 
185  m_CommandQue = resources->GetCommandQueue();
186 
187  cl_int clErr = 0;
188  m_Initialized = CHECK_OCL_ERR(clErr);
189 
190  if ( m_ClFiles.empty())
191  {
192  MITK_ERROR<<"No OpenCL Source FILE specified";
193  return false;
194  }
195 
196  if (m_ClProgram == nullptr)
197  {
198  try
199  {
200  this->m_ClProgram = resources->GetProgram( this->m_FilterID );
201  }
202  catch(const mitk::Exception& e)
203  {
204  MITK_INFO << "Program not stored in resource manager, compiling. " << e;
205  this->CompileSource();
206  }
207  }
208 
209  return m_Initialized;
210 }
211 
212 void mitk::OclFilter::LoadSourceFiles(CStringList &sourceCode, ClSizeList &sourceCodeSize)
213 {
214  for( CStringList::iterator it = m_ClFiles.begin(); it != m_ClFiles.end(); ++it )
215  {
216  MITK_DEBUG << "Load file :" << *it;
218 
219  if( !mdr.IsValid() )
220  MITK_WARN << "Could not load resource: " << mdr.GetName() << " is invalid!";
221 
222  us::ModuleResourceStream rss(mdr);
223 
224  // read resource file to a string
225  std::istreambuf_iterator<char> eos;
226  std::string source(std::istreambuf_iterator<char>(rss), eos);
227 
228  // add preambel and build up string to compile
229  std::string src(m_Preambel);
230  src.append("\n");
231  src.append(source);
232 
233  // allocate new char buffer
234  char* tmp = new char[src.size() + 1];
235  strcpy(tmp,src.c_str());
236 
237  // add source to list
238  sourceCode.push_back((const char*)tmp);
239  sourceCodeSize.push_back(src.size());
240  }
241 }
242 
244 {
245  // helper variable
246  int clErr = 0;
247  CStringList sourceCode;
248  ClSizeList sourceCodeSize;
249 
250  if (m_ClFiles.empty())
251  {
252  MITK_ERROR("ocl.filter") << "No shader source file was set";
253  return;
254  }
255 
256  //get a valid opencl context
259 
260  cl_context gpuContext = resources->GetContext();
261  // load the program source from file
262  LoadSourceFiles(sourceCode, sourceCodeSize);
263 
264  if ( !sourceCode.empty() )
265  {
266  // create program from all files in the file list
267  m_ClProgram = clCreateProgramWithSource(gpuContext, sourceCode.size(), &sourceCode[0], &sourceCodeSize[0], &clErr);
268  CHECK_OCL_ERR(clErr);
269 
270  // build the source code
271  MITK_DEBUG << "Building Program Source";
272  std::string compilerOptions = "";
273  compilerOptions.append(m_ClCompilerFlags);
274 
275  MITK_DEBUG("ocl.filter") << "cl compiler flags: " << compilerOptions.c_str();
276 
277  clErr = clBuildProgram(m_ClProgram, 0, nullptr, compilerOptions.c_str(), nullptr, nullptr);
278  CHECK_OCL_ERR(clErr);
279 
280  // if OpenCL Source build failed
281  if (clErr != CL_SUCCESS)
282  {
283  MITK_ERROR("ocl.filter") << "Failed to build source";
284  oclLogBuildInfo(m_ClProgram, resources->GetCurrentDevice() );
285  oclLogBinary(m_ClProgram, resources->GetCurrentDevice() );
286  m_Initialized = false;
287  }
288 
289  // store the succesfully build program into the program storage provided by the resource service
290  resources->InsertProgram(m_ClProgram, m_FilterID, true);
291 
292  // free the char buffers with the source code
293  for( CStringList::iterator it = sourceCode.begin(); it != sourceCode.end(); ++it )
294  {
295  delete[] *it;
296  }
297  }
298  else
299  {
300  MITK_ERROR("ocl.filter") << "Could not load from source";
301  m_Initialized = false;
302  }
303 }
304 
305 void mitk::OclFilter::SetWorkingSize(unsigned int locx, unsigned int dimx, unsigned int locy, unsigned int dimy, unsigned int locz, unsigned int dimz)
306 {
307  // set the local work size
308  this->m_LocalWorkSize[0] = locx;
309  this->m_LocalWorkSize[1] = locy;
310  this->m_LocalWorkSize[2] = locz;
311 
312  this->m_GlobalWorkSize[0] = dimx;
313  this->m_GlobalWorkSize[1] = dimy;
314  this->m_GlobalWorkSize[2] = dimz;
315 
316  // estimate the global work size
317  this->m_GlobalWorkSize[0] = iDivUp( dimx, this->m_LocalWorkSize[0]) * this->m_LocalWorkSize[0];
318 
319  if ( dimy > 1)
320  this->m_GlobalWorkSize[1] = iDivUp( dimy, this->m_LocalWorkSize[1]) * this->m_LocalWorkSize[1];
321  if( dimz > 1 )
322  this->m_GlobalWorkSize[2] = iDivUp( dimz, this->m_LocalWorkSize[2]) * this->m_LocalWorkSize[2];
323 }
324 
325 void mitk::OclFilter::SetSourcePreambel(const char* preambel)
326 {
327  this->m_Preambel = preambel;
328 }
329 
330 void mitk::OclFilter::AddSourceFile(const char* filename)
331 {
332  m_ClFiles.push_back(filename);
333 }
334 
335 void mitk::OclFilter::SetCompilerFlags(const char* flags)
336 {
337  m_ClCompilerFlags = flags;
338 }
339 
340 
342 {
343  return m_Initialized;
344 }
345 
347 {
349  auto device = resources->GetCurrentDevice();
350  return oclGetGlobalMemSize(device);
351 }
#define CHECK_OCL_ERR(_er)
Definition: mitkOclUtils.h:21
unsigned int iDivUp(unsigned int dividend, unsigned int divisor)
Method to estimate an integer quotient C from given dividend and divisor higher or equal to the corre...
ServiceReferenceU GetServiceReference(const std::string &clazz)
std::string m_FilterID
Unique ID of the filter, needs to be specified in the constructor of the derived class.
Definition: mitkOclFilter.h:86
OclFilter()
Constructor.
#define MITK_INFO
Definition: mitkLogMacros.h:18
bool m_Initialized
status of the filter
Definition: mitkOclFilter.h:95
void SetSourcePreambel(const char *preambel)
Add some source code on the beginning of the loaded source.
cl_program m_ClProgram
The compiled OpenCL program.
Definition: mitkOclFilter.h:80
#define MITK_ERROR
Definition: mitkLogMacros.h:20
const char * m_Preambel
source preambel for e.g. #define commands to be inserted into the OpenCL source
Definition: mitkOclFilter.h:89
bool Initialize()
Initialize all necessary parts of the filter.
#define MITK_DEBUG
Definition: mitkLogMacros.h:22
vcl_size_t m_GlobalWorkSize[3]
The global work size of the filter.
void oclLogBinary(cl_program clProg, cl_device_id clDev)
Logs the GPU Program binary code.
void CompileSource()
Compile the program source.
virtual bool IsInitialized()
Returns true if the initialization was successfull.
bool ExecuteKernelChunksInBatches(cl_kernel kernel, unsigned int workSizeDim, vcl_size_t *chunksDim, vcl_size_t batchSize, int waitTimems)
Execute the given kernel on the OpenCL Index-Space defined by the local and global work sizes...
bool ExecuteKernelChunks(cl_kernel kernel, unsigned int workSizeDim, vcl_size_t *chunksDim)
Execute the given kernel on the OpenCL Index-Space defined by the local and global work sizes...
std::vector< const char * > CStringList
Definition: mitkOclFilter.h:67
virtual unsigned long GetDeviceMemory()
Returns the amount of global memory of the used device in bytes.
void * GetService(const ServiceReferenceBase &reference)
static Vector3D offset
virtual ~OclFilter()
Destructor.
#define MITK_WARN
Definition: mitkLogMacros.h:19
An object of this class represents an exception of MITK. Please don&#39;t instantiate exceptions manually...
Definition: mitkException.h:45
std::vector< vcl_size_t > ClSizeList
Definition: mitkOclFilter.h:68
ModuleResource GetResource(const std::string &path) const
Definition: usModule.cpp:267
CStringList m_ClFiles
List of sourcefiles that will be compiled for this filter.
Definition: mitkOclFilter.h:92
void oclLogBuildInfo(cl_program clProg, cl_device_id clDev)
Shows the OpenCL-Program build info, called if clBuildProgram != CL_SUCCES.
virtual void RemoveProgram(const std::string &name)=0
Remove given program from storage.
void SetCompilerFlags(const char *flags)
Set specific compilerflags to compile the CL source. Default is set to nullptr; example: "-cl-fast-re...
cl_command_queue m_CommandQue
Command queue for the filter.
Definition: mitkOclFilter.h:83
virtual us::Module * GetModule()=0
Get the Module of the filter. Needs to be implemented by every subclass. The filter will load the Ope...
cl_ulong oclGetGlobalMemSize(cl_device_id device)
Returns the Global memory size of the current device.
vcl_size_t m_LocalWorkSize[3]
The local work size fo the filter.
Definition: mitkOclFilter.h:98
virtual cl_context GetContext() const =0
Returns a valid OpenCL Context (if applicable) or nullptr if none present.
void SetWorkingSize(unsigned int locx, unsigned int dimx, unsigned int locy=1, unsigned int dimy=1, unsigned int locz=1, unsigned int dimz=1)
Set the working size for the following OpenCL kernel call.
bool ExecuteKernel(cl_kernel kernel, unsigned int workSizeDim)
Execute the given kernel on the OpenCL Index-Space defined by the local and global work sizes...
void LoadSourceFiles(CStringList &SourceCodeList, ClSizeList &SourceCodeSizeList)
Helper functions that load sourcefiles from the module context in the Initialize function.
const char * m_ClCompilerFlags
String that contains the compiler flags.
Definition: mitkOclFilter.h:77
static ModuleContext * GetModuleContext()
Returns the module context of the calling module.
virtual cl_device_id GetCurrentDevice() const =0
Returns the identifier of an OpenCL device related to the current context.
void AddSourceFile(const char *filename)
Add a source file from the resource files to the OpenCL shader file list. Multiple files can be added...
std::string GetName() const
Declaration of the OpenCL Resources micro-service.