diff --git a/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.cpp b/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.cpp index b44c7c6463..4e7dede7a7 100644 --- a/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.cpp +++ b/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.cpp @@ -1,97 +1,100 @@ /*=================================================================== The Medical Imaging Interaction Toolkit (MITK) Copyright (c) German Cancer Research Center, Division of Medical and Biological Informatics. All rights reserved. This software is distributed WITHOUT ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See LICENSE.txt or http://www.mitk.org for details. ===================================================================*/ #include "mitkOclBinaryThresholdImageFilter.h" mitk::OclBinaryThresholdImageFilter::OclBinaryThresholdImageFilter() : m_ckBinaryThreshold( NULL ) { std::string path = "BinaryThresholdFilter.cl"; this->SetSourceFile( path.c_str() ); this->m_FilterID = "BinaryThreshold"; this->m_LowerThreshold = 10; this->m_UpperThreshold = 200; this->m_InsideValue = 100; this->m_OutsideValue = 0; } mitk::OclBinaryThresholdImageFilter::~OclBinaryThresholdImageFilter() { if ( this->m_ckBinaryThreshold ) { clReleaseKernel( m_ckBinaryThreshold ); } } void mitk::OclBinaryThresholdImageFilter::Update() { //Check if context & program available if (!this->Initialize()) { + mitk::ServiceReference ref = GetModuleContext()->GetServiceReference(); + OclResourceService* resources = GetModuleContext()->GetService(ref); + // clean-up also the resources - OpenCLActivator::GetResourceServiceRef()->InvalidateStorage(); + resources->InvalidateStorage(); mitkThrow() <<"Filter is not initialized. Cannot update."; } else{ // Execute this->Execute(); } } void mitk::OclBinaryThresholdImageFilter::Execute() { cl_int clErr = 0; try { this->InitExec( this->m_ckBinaryThreshold ); } catch( const mitk::Exception& e) { MITK_ERROR << "Catched exception while initializing filter: " << e.what(); return; } // set kernel arguments clErr = clSetKernelArg( this->m_ckBinaryThreshold, 2, sizeof(cl_int), &(this->m_LowerThreshold) ); clErr |= clSetKernelArg( this->m_ckBinaryThreshold, 3, sizeof(cl_int), &(this->m_UpperThreshold) ); clErr |= clSetKernelArg( this->m_ckBinaryThreshold, 4, sizeof(cl_int), &(this->m_OutsideValue) ); clErr |= clSetKernelArg( this->m_ckBinaryThreshold, 5, sizeof(cl_int), &(this->m_InsideValue) ); CHECK_OCL_ERR( clErr ); // execute the filter on a 3D NDRange this->ExecuteKernel( m_ckBinaryThreshold, 3); // signalize the GPU-side data changed m_Output->Modified( GPU_DATA ); } bool mitk::OclBinaryThresholdImageFilter::Initialize() { bool buildErr = true; cl_int clErr = 0; if ( OclFilter::Initialize() ) { this->m_ckBinaryThreshold = clCreateKernel( this->m_ClProgram, "ckBinaryThreshold", &clErr); buildErr |= CHECK_OCL_ERR( clErr ); } return (Superclass::IsInitialized() && buildErr ); } diff --git a/Modules/OpenCL/mitkOclFilter.cpp b/Modules/OpenCL/mitkOclFilter.cpp index ea91e778bf..4321af78bf 100644 --- a/Modules/OpenCL/mitkOclFilter.cpp +++ b/Modules/OpenCL/mitkOclFilter.cpp @@ -1,219 +1,228 @@ /*=================================================================== The Medical Imaging Interaction Toolkit (MITK) Copyright (c) German Cancer Research Center, Division of Medical and Biological Informatics. All rights reserved. This software is distributed WITHOUT ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See LICENSE.txt or http://www.mitk.org for details. ===================================================================*/ //Ocl #include "mitkOclFilter.h" #include "mitkOclUtils.h" #include "mitkOpenCLActivator.h" //Mitk #include #include #include #include mitk::OclFilter::OclFilter() : m_ClFile(), m_ClSource(NULL), m_ClCompilerFlags(""), m_ClProgram(NULL), m_CommandQue(NULL), m_FilterID("mitkOclFilter"), m_Preambel(" "), m_Initialized(false) { m_ClSourcePath = MITK_ROOT; m_ClSourcePath += "Modules/OpenCL/ShaderSources"; } mitk::OclFilter::OclFilter(const char* filename) : m_ClFile(), m_ClSource(NULL), m_ClCompilerFlags(""), m_ClProgram(NULL), m_CommandQue(NULL), m_FilterID(filename), m_Preambel(" "), m_Initialized(false) { m_ClSourcePath = MITK_ROOT; m_ClSourcePath += "Modules/OpenCL/ShaderSources"; } mitk::OclFilter::~OclFilter() { MITK_DEBUG << "OclFilter Destructor"; // release program if (m_ClProgram) { cl_int clErr = 0; + mitk::ServiceReference ref = GetModuleContext()->GetServiceReference(); + OclResourceService* resources = GetModuleContext()->GetService(ref); + // remove program from storage - OpenCLActivator::GetResourceServiceRef()->RemoveProgram(m_FilterID); + resources->RemoveProgram(m_FilterID); // release program clErr = clReleaseProgram(this->m_ClProgram); CHECK_OCL_ERR(clErr); } } bool mitk::OclFilter::ExecuteKernel( cl_kernel kernel, unsigned int workSizeDim ) { cl_int clErr = 0; clErr = clEnqueueNDRangeKernel( this->m_CommandQue, kernel, workSizeDim, NULL, this->m_GlobalWorkSize, m_LocalWorkSize, 0, NULL, NULL); CHECK_OCL_ERR( clErr ); return ( clErr == CL_SUCCESS ); } bool mitk::OclFilter::Initialize() { - m_CommandQue = OpenCLActivator::GetResourceServiceRef()->GetCommandQueue(); + mitk::ServiceReference ref = GetModuleContext()->GetServiceReference(); + OclResourceService* resources = GetModuleContext()->GetService(ref); + + m_CommandQue = resources->GetCommandQueue(); cl_int clErr = 0; m_Initialized = CHECK_OCL_ERR(clErr); if ((m_ClSource==NULL) && (m_ClFile.empty())) { MITK_ERROR<<"No OpenCL Source FILE specified"; return false; } if (m_ClProgram == NULL) { try { - this->m_ClProgram = OpenCLActivator::GetResourceServiceRef()->GetProgram( this->m_FilterID ); + this->m_ClProgram = resources->GetProgram( this->m_FilterID ); } catch(const mitk::Exception& e) { MITK_INFO << "Program not stored in resource manager, compiling."; this->CompileSource(); } } return m_Initialized; } void mitk::OclFilter::SetSourceFile(const char* filename) { MITK_DEBUG("ocl.filter") << "Setting source [" << filename <<" ]"; mitk::StandardFileLocations::GetInstance()->AddDirectoryForSearch( m_ClSourcePath.c_str(), true); // search for file m_ClFile = mitk::StandardFileLocations::GetInstance()->FindFile( filename); } void mitk::OclFilter::CompileSource() { if (m_ClFile.empty() && m_ClSource == NULL) { MITK_ERROR("ocl.filter") << "No shader source file was set"; return; } // help variables size_t szKernelLength; int clErr = 0; //get a valid opencl context - cl_context gpuContext = OpenCLActivator::GetResourceServiceRef()->GetContext(); + mitk::ServiceReference ref = GetModuleContext()->GetServiceReference(); + OclResourceService* resources = GetModuleContext()->GetService(ref); + + cl_context gpuContext = resources->GetContext(); // load the program source from file m_ClSource = oclLoadProgramSource( m_ClFile.c_str(), this->m_Preambel, &szKernelLength); if (m_ClSource != NULL) { m_ClProgram = clCreateProgramWithSource(gpuContext, 1, (const char**)&m_ClSource, &szKernelLength, &clErr); CHECK_OCL_ERR(clErr); // build the source code MITK_DEBUG << "Building Program Source"; std::string compilerOptions = ""; compilerOptions.append(m_ClCompilerFlags); // activate the include compiler flag compilerOptions.append(" -I"); // set the path of the current gpu source dir as opencl // include folder compilerOptions.append(m_ClSourcePath.c_str()); MITK_DEBUG("ocl.filter") << "cl compiler flags: " << compilerOptions.c_str(); clErr = clBuildProgram(m_ClProgram, 0, NULL, compilerOptions.c_str(), NULL, NULL); CHECK_OCL_ERR(clErr); // if OpenCL Source build failed if (clErr != CL_SUCCESS) { MITK_ERROR("ocl.filter") << "Failed to build source"; - oclLogBuildInfo(m_ClProgram, OpenCLActivator::GetResourceServiceRef()->GetCurrentDevice() ); - oclLogBinary(m_ClProgram, OpenCLActivator::GetResourceServiceRef()->GetCurrentDevice() ); + oclLogBuildInfo(m_ClProgram, resources->GetCurrentDevice() ); + oclLogBinary(m_ClProgram, resources->GetCurrentDevice() ); m_Initialized = false; } // store the succesfully build program into the program storage provided by the resource service - OpenCLActivator::GetResourceServiceRef()->InsertProgram(m_ClProgram, m_FilterID, true); + resources->InsertProgram(m_ClProgram, m_FilterID, true); } else { MITK_ERROR("ocl.filter") << "Could not load from source"; m_Initialized = false; } } void mitk::OclFilter::SetWorkingSize(unsigned int locx, unsigned int dimx, unsigned int locy, unsigned int dimy, unsigned int locz, unsigned int dimz) { // set the local work size this->m_LocalWorkSize[0] = locx; this->m_LocalWorkSize[1] = locy; this->m_LocalWorkSize[2] = locz; // estimate the global work size this->m_GlobalWorkSize[0] = iDivUp( dimx, this->m_LocalWorkSize[0]) *this->m_LocalWorkSize[0]; this->m_GlobalWorkSize[1] = iDivUp( dimy, this->m_LocalWorkSize[1]) * this->m_LocalWorkSize[1]; if( dimz <= 1 ) this->m_GlobalWorkSize[2] = 1; else this->m_GlobalWorkSize[2] = iDivUp( dimz, this->m_LocalWorkSize[2]) * this->m_LocalWorkSize[2]; } void mitk::OclFilter::SetSourcePreambel(const char* preambel) { this->m_Preambel = preambel; } void mitk::OclFilter::SetSourcePath(const char* path) { m_ClSourcePath = path; } void mitk::OclFilter::SetCompilerFlags(const char* flags) { m_ClCompilerFlags = flags; } bool mitk::OclFilter::IsInitialized() { return m_Initialized; } diff --git a/Modules/OpenCL/mitkOclImage.cpp b/Modules/OpenCL/mitkOclImage.cpp index 4bf61e72b7..dd1a8d4e12 100644 --- a/Modules/OpenCL/mitkOclImage.cpp +++ b/Modules/OpenCL/mitkOclImage.cpp @@ -1,347 +1,354 @@ /*=================================================================== The Medical Imaging Interaction Toolkit (MITK) Copyright (c) German Cancer Research Center, Division of Medical and Biological Informatics. All rights reserved. This software is distributed WITHOUT ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See LICENSE.txt or http://www.mitk.org for details. ===================================================================*/ #include "mitkOclImage.h" #include "mitkImageDataItem.h" #include "mitkCommon.h" #include "mitkLogMacros.h" #include "mitkOclUtils.h" #include mitk::OclImage::OclImage() : m_gpuImage(NULL), m_context(NULL), m_bufferSize(0), m_gpuModified(false), m_cpuModified(false), m_Image(NULL), m_dim(0), m_Dims(NULL), m_BpE(1), m_formatSupported(false) { } mitk::OclImage::~OclImage() { MITK_INFO << "OclImage Destructor"; //release GMEM Image buffer if (m_gpuImage) clReleaseMemObject(m_gpuImage); } cl_mem mitk::OclImage::CreateGPUImage(unsigned int _wi, unsigned int _he, unsigned int _de, unsigned int _bpp) { MITK_INFO << "CreateGPUImage call with: BPP=" << _bpp; this->m_Dims = new unsigned int[MAX_DIMS]; m_Dims[0] = _wi; m_Dims[1] = _he; m_Dims[2] = _de; for (unsigned int i=3; iGetContext(); + mitk::ServiceReference ref = GetModuleContext()->GetServiceReference(); + OclResourceService* resources = GetModuleContext()->GetService(ref); + + cl_context gpuContext = resources->GetContext(); int clErr; m_gpuImage = clCreateBuffer( gpuContext, CL_MEM_READ_WRITE, m_bufferSize * m_BpE, NULL, &clErr); CHECK_OCL_ERR(clErr); return m_gpuImage; } void mitk::OclImage::InitializeByMitkImage(mitk::Image::Pointer _image) { this->m_Image = _image; this->m_cpuModified = true; this->m_gpuModified = false; this->m_gpuImage = NULL; // compute the size of the GMEM buffer this->m_dim = this->m_Image->GetDimension(); this->m_Dims = this->m_Image->GetDimensions(); MITK_INFO << "Image: " << this->m_Dims[0] <<"x"<< this->m_Dims[1] <<"x"<< this->m_Dims[2]; // get the dimensions this->m_bufferSize = 1; for (unsigned int i=0; im_dim; i++) { this->m_bufferSize *= this->m_Dims[i]; } // multiply by sizeof(PixelType) this->m_BpE = ( this->m_Image->GetPixelType().GetBpe() / 8); } bool mitk::OclImage::IsModified(int _type) { if (_type) return m_cpuModified; else return m_gpuModified; } void mitk::OclImage::Modified(int _type) { // defines... GPU: 0, CPU: 1 m_cpuModified = _type; m_gpuModified = !_type; } int mitk::OclImage::TransferDataToGPU(cl_command_queue gpuComQueue) { cl_int clErr = 0; // check whether an image present if (!m_Image->IsInitialized()){ MITK_ERROR("ocl.Image") << "(mitk) Image not initialized!\n"; return -1; } // there is a need for copy only if RAM-Data newer then GMEM data if (m_cpuModified) { //check the buffer if(m_gpuImage == NULL) { clErr = this->AllocateGPUImage(); } if (m_Image->IsInitialized() && (clErr == CL_SUCCESS)) { const size_t origin[3] = {0, 0, 0}; const size_t region[3] = {m_Dims[0], m_Dims[1], m_Dims[2]}; if( this->m_formatSupported ) { clErr = clEnqueueWriteImage( gpuComQueue, m_gpuImage, CL_TRUE, origin, region, 0, 0, m_Image->GetData(), 0, NULL, NULL); } else { MITK_ERROR << "Selected image format currently not supported..."; } CHECK_OCL_ERR(clErr); } m_gpuModified = true; } return clErr; } cl_int mitk::OclImage::AllocateGPUImage() { cl_int clErr = 0; - cl_context gpuContext =OpenCLActivator::GetResourceServiceRef()->GetContext(); + + mitk::ServiceReference ref = GetModuleContext()->GetServiceReference(); + OclResourceService* resources = GetModuleContext()->GetService(ref); + + cl_context gpuContext = resources->GetContext(); // initialize both proposed and supported format variables to same value this->m_proposedFormat = this->ConvertPixelTypeToOCLFormat(); this->m_supportedFormat = this->m_proposedFormat; // test the current format for HW support - this->m_formatSupported = OpenCLActivator::GetResourceServiceRef()->GetIsFormatSupported( &(this->m_supportedFormat) ); + this->m_formatSupported = resources->GetIsFormatSupported( &(this->m_supportedFormat) ); // create an transfer kernel object in case the proposed format is not supported if( !(this->m_formatSupported) ) { MITK_ERROR << "Original format not supported on the installed graphics card."; return -1; } // create new buffer if( this->m_dim > 2) { //Create a 3D Image m_gpuImage = clCreateImage3D(gpuContext, CL_MEM_READ_ONLY, &m_supportedFormat, *(m_Dims), *(m_Dims+1), *(m_Dims+2), 0, 0, NULL, &clErr); } else { //Create a 2D Image m_gpuImage = clCreateImage2D(gpuContext, CL_MEM_READ_ONLY, &m_supportedFormat, *(m_Dims), *(m_Dims+1), 0, NULL, &clErr); } CHECK_OCL_ERR(clErr); return clErr; } cl_mem mitk::OclImage::GetGPUImage(cl_command_queue gpuComQueue) { // clGetMemObjectInfo() cl_mem_object_type memInfo; cl_int clErr = 0; // query image object info only if already initialized if( this->m_gpuImage ) { clErr = clGetMemObjectInfo(this->m_gpuImage, CL_MEM_TYPE, sizeof(cl_mem_object_type), &memInfo, NULL ); CHECK_OCL_ERR(clErr); } MITK_INFO << "Querying info for object, recieving: " << memInfo; // test if m_gpuImage CL_MEM_IMAGE_2/3D // if not, copy buffer to image if (memInfo == CL_MEM_OBJECT_BUFFER) { MITK_WARN << " Passed oclImage is a buffer-object, creating image"; //hold a copy of the buffer gmem pointer cl_mem tempBuffer = this->m_gpuImage; const size_t origin[3] = {0, 0, 0}; size_t region[3] = {this->m_Dims[0], this->m_Dims[1], 1}; clErr = this->AllocateGPUImage(); this->m_dim = 3; //copy last data to the image data clErr = clEnqueueCopyBufferToImage( gpuComQueue, tempBuffer, m_gpuImage, 0, origin, region, 0, NULL, NULL); CHECK_OCL_ERR(clErr); //release pointer clReleaseMemObject(tempBuffer); } return m_gpuImage; } void mitk::OclImage::SetPixelType(const cl_image_format *_image) { this->m_proposedFormat.image_channel_data_type = _image->image_channel_data_type; this->m_proposedFormat.image_channel_order = _image->image_channel_order; } void* mitk::OclImage::TransferDataToCPU(cl_command_queue gpuComQueue) { cl_int clErr = 0; // if image created on GPU, needs to create mitk::Image if( m_Image.IsNull() ){ MITK_INFO << "Image not initialized, creating new one."; m_Image = mitk::Image::New(); } // check buffersize/image size char* data = new char[m_bufferSize * m_BpE]; // debug info oclPrintMemObjectInfo( m_gpuImage ); clErr = clEnqueueReadBuffer( gpuComQueue, m_gpuImage, CL_FALSE, 0, m_bufferSize * m_BpE, data ,0, NULL, NULL); CHECK_OCL_ERR(clErr); clFlush( gpuComQueue ); // the cpu data is same as gpu this->m_gpuModified = false; return (void*) data; } cl_image_format mitk::OclImage::ConvertPixelTypeToOCLFormat() { cl_image_format texFormat; //single channel Gray-Valued Images texFormat.image_channel_order = CL_R; MITK_INFO << "Class own value: " << this->m_BpE; switch ( this->m_BpE ) { case 1: texFormat.image_channel_data_type = CL_UNSIGNED_INT8; MITK_INFO<< "PixelType: UCHAR => CLFormat: [CL_UNORM_INT8, CL_R]"; break; case 2: texFormat.image_channel_data_type = CL_SIGNED_INT16; // texFormat.image_channel_order = CL_R; MITK_INFO<< "PixelType: SHORT => CLFormat: [CL_SIGNED_INT16, CL_R]"; break; case 4: texFormat.image_channel_data_type = CL_FLOAT; MITK_INFO<< "Choosing CL_FLOAT"; break; default: texFormat.image_channel_data_type = CL_UNORM_INT8; texFormat.image_channel_order = CL_RG; MITK_INFO<< "Choosing (default) short: 2-Channel UCHAR"; break; } return texFormat; } int mitk::OclImage::GetDimension(int idx) const { if (this->m_dim > idx) { return m_Dims[idx]; } else { MITK_WARN << "Trying to access non-existing dimension."; return 1; } } void mitk::OclImage::SetDimensions(unsigned int* Dims) { m_Dims = Dims; } void mitk::OclImage::SetDimension(unsigned short dim) { m_dim = dim; } float mitk::OclImage::GetSpacing(int idx) { if (this->m_dim > idx) { const float* imSpacing = m_Image->GetSlicedGeometry()->GetFloatSpacing(); return imSpacing[idx]; } else { MITK_WARN << "Trying to access non-existing dimension."; return 1; } } void mitk::OclImage::InitializeMITKImage() { this->m_Image = mitk::Image::New(); } void mitk::OclImage::GetOffset(float* _imOffset) const { itk::Vector result2; result2.Fill(0.0f); result2 = this->m_Image->GetGeometry()->GetIndexToWorldTransform()->GetOffset(); _imOffset[0] = result2[0]; _imOffset[1] = result2[1]; _imOffset[2] = result2[2]; } diff --git a/Modules/OpenCL/mitkOpenCLActivator.cpp b/Modules/OpenCL/mitkOpenCLActivator.cpp index 28e7d6fff4..29e2efb2c0 100644 --- a/Modules/OpenCL/mitkOpenCLActivator.cpp +++ b/Modules/OpenCL/mitkOpenCLActivator.cpp @@ -1,42 +1,33 @@ /*=================================================================== The Medical Imaging Interaction Toolkit (MITK) Copyright (c) German Cancer Research Center, Division of Medical and Biological Informatics. All rights reserved. This software is distributed WITHOUT ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See LICENSE.txt or http://www.mitk.org for details. ===================================================================*/ #include "mitkOpenCLActivator.h" void OpenCLActivator::Load(mitk::ModuleContext *context) { // generate context m_ResourceService.reset(new OclResourceServiceImpl); mitk::ServiceProperties props; context->RegisterService(m_ResourceService.get(), props); - m_InternalResourceReference = m_ResourceService.get(); } void OpenCLActivator::Unload(mitk::ModuleContext *) { - m_InternalResourceReference = NULL; m_ResourceService.release(); } -OclResourceService* OpenCLActivator::GetResourceServiceRef() -{ - return m_InternalResourceReference; -} - -OclResourceServiceImpl* OpenCLActivator::m_InternalResourceReference = NULL; - US_EXPORT_MODULE_ACTIVATOR(MitkOcl, OpenCLActivator ) diff --git a/Modules/OpenCL/mitkOpenCLActivator.h b/Modules/OpenCL/mitkOpenCLActivator.h index fcc88a3b9f..c3132404d0 100644 --- a/Modules/OpenCL/mitkOpenCLActivator.h +++ b/Modules/OpenCL/mitkOpenCLActivator.h @@ -1,59 +1,54 @@ /*=================================================================== The Medical Imaging Interaction Toolkit (MITK) Copyright (c) German Cancer Research Center, Division of Medical and Biological Informatics. All rights reserved. This software is distributed WITHOUT ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See LICENSE.txt or http://www.mitk.org for details. ===================================================================*/ #ifndef __mitkOpenCLActivator_h #define __mitkOpenCLActivator_h #include "mitkOclResourceServiceImpl_p.h" #include #include #include #include #include #include #include /** * @class OpenCLActivator * * @brief Custom activator for the OpenCL Module in order to register * and provide the OclResourceService */ class US_ABI_LOCAL OpenCLActivator : public mitk::ModuleActivator { private: std::auto_ptr m_ResourceService; - /** Static variable for fast internal access to the ResourceService */ - static OclResourceServiceImpl* m_InternalResourceReference; - public: /** @brief Load module context */ void Load(mitk::ModuleContext *context); /** @brief Unload module context */ void Unload(mitk::ModuleContext* ); - static OclResourceService* GetResourceServiceRef(); - }; #endif // __mitkOpenCLActivator_h