diff --git a/Modules/OpenCL/Testing/mitkOclBinaryThresholdImageFilterTest.cpp b/Modules/OpenCL/Testing/mitkOclBinaryThresholdImageFilterTest.cpp index cb57f9babe..f92fc78268 100644 --- a/Modules/OpenCL/Testing/mitkOclBinaryThresholdImageFilterTest.cpp +++ b/Modules/OpenCL/Testing/mitkOclBinaryThresholdImageFilterTest.cpp @@ -1,111 +1,111 @@ /*=================================================================== 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 #include #include #include // itk filter for reference computation #include #include #include using namespace mitk; /** This function is testing the class mitk::OclContextManager. */ int mitkOclBinaryThresholdImageFilterTest( int argc, char* argv[] ) { MITK_TEST_BEGIN("mitkOclBinaryThresholdImageFilterTest"); - ServiceReference ref = GetModuleContext()->GetServiceReference(); + us::ServiceReference ref = GetModuleContext()->GetServiceReference(); MITK_TEST_CONDITION_REQUIRED( ref != 0, "Got valid ServiceReference" ); OclResourceService* resources = GetModuleContext()->GetService(ref); MITK_TEST_CONDITION_REQUIRED( resources != NULL, "OpenCL Resource service available." ); cl_context gpuContext = resources->GetContext(); MITK_TEST_CONDITION_REQUIRED( gpuContext != NULL, "Got not-null OpenCL context."); cl_device_id gpuDevice = resources->GetCurrentDevice(); MITK_TEST_CONDITION_REQUIRED( gpuDevice != NULL, "Got not-null OpenCL device."); //Create a random reference image mitk::Image::Pointer inputImage = mitk::ImageGenerator::GenerateRandomImage(119, 204, 52, 1, // dimension 1.0f, 1.0f, 1.0f, // spacing 255, 0); // max, min MITK_TEST_CONDITION_REQUIRED( inputImage.IsNotNull(), "Input (random) mitk::Image object instantiated."); // FIXME: could also be random values int upperThr = 255; int lowerThr = 60; int outsideVal = 0; int insideVal = 100; mitk::OclBinaryThresholdImageFilter* oclFilter = new mitk::OclBinaryThresholdImageFilter; MITK_TEST_CONDITION_REQUIRED( oclFilter != NULL, "Filter was created. "); oclFilter->SetInput( inputImage ); oclFilter->SetUpperThreshold( upperThr ); oclFilter->SetLowerThreshold( lowerThr ); oclFilter->SetOutsideValue( outsideVal ); oclFilter->SetInsideValue( insideVal ); oclFilter->Update(); mitk::Image::Pointer outputImage = mitk::Image::New(); outputImage = oclFilter->GetOutput(); MITK_TEST_CONDITION_REQUIRED( outputImage.IsNotNull(), "Filter returned an not-NULL image. "); // reference computation typedef itk::Image< unsigned char, 3> ImageType; typedef itk::BinaryThresholdImageFilter< ImageType, ImageType > ThresholdFilterType; ImageType::Pointer itkInputImage = ImageType::New(); CastToItkImage( inputImage, itkInputImage ); ThresholdFilterType::Pointer refThrFilter = ThresholdFilterType::New(); refThrFilter->SetInput( itkInputImage ); refThrFilter->SetLowerThreshold( lowerThr ); refThrFilter->SetUpperThreshold( upperThr ); refThrFilter->SetOutsideValue( outsideVal ); refThrFilter->SetInsideValue( insideVal ); typedef itk::SubtractImageFilter< ImageType, ImageType > SubtractFilterType; SubtractFilterType::Pointer subFilt = SubtractFilterType::New(); ImageType::Pointer gpuReferenceImage = ImageType::New(); CastToItkImage( oclFilter->GetOutput() ,gpuReferenceImage ); subFilt->SetInput1( refThrFilter->GetOutput() ); subFilt->SetInput2( gpuReferenceImage ); typedef itk::StatisticsImageFilter< ImageType > StatFilterType; StatFilterType::Pointer stats = StatFilterType::New(); stats->SetInput( subFilt->GetOutput() ); stats->Update(); MITK_TEST_CONDITION( stats->GetMaximum() == 0, "Maximal value in the difference image is 0."); MITK_TEST_CONDITION( stats->GetMinimum() == 0, "Minimal value in the difference image is 0.") MITK_TEST_END(); -} +} \ No newline at end of file diff --git a/Modules/OpenCL/Testing/mitkOclImageTest.cpp b/Modules/OpenCL/Testing/mitkOclImageTest.cpp index 38ae3b1a62..1474300709 100644 --- a/Modules/OpenCL/Testing/mitkOclImageTest.cpp +++ b/Modules/OpenCL/Testing/mitkOclImageTest.cpp @@ -1,94 +1,94 @@ /*=================================================================== 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 #include #include using namespace mitk; /** This function is testing the mitk::OclImage class. */ int mitkOclImageTest( int /*argc*/, char* /*argv*/[] ) { MITK_TEST_BEGIN("mitkOclImageTest"); - ServiceReference ref = GetModuleContext()->GetServiceReference(); + us::ServiceReference ref = GetModuleContext()->GetServiceReference(); MITK_TEST_CONDITION_REQUIRED( ref != 0, "Got valid ServiceReference" ); OclResourceService* resources = GetModuleContext()->GetService(ref); MITK_TEST_CONDITION_REQUIRED( resources != NULL, "OpenCL Resource service available." ); cl_context gpuContext = resources->GetContext(); MITK_TEST_CONDITION_REQUIRED( gpuContext != NULL, "Got not-null OpenCL context."); cl_device_id gpuDevice = resources->GetCurrentDevice(); MITK_TEST_CONDITION_REQUIRED( gpuDevice != NULL, "Got not-null OpenCL device."); //Create a random reference image mitk::Image::Pointer reference = mitk::ImageGenerator::GenerateRandomImage(119, 204, 52, 1, // dimension 1.0f, 1.0f, 1.0f, // spacing 1024, 0); // max, min MITK_TEST_CONDITION_REQUIRED( reference.IsNotNull(), "Reference mitk::Image object instantiated."); mitk::OclImage::Pointer first = mitk::OclImage::New(); first->InitializeByMitkImage(reference); MITK_TEST_CONDITION_REQUIRED(first.IsNotNull(), "oclImage object instantiated."); // test if oclImage correct initialized MITK_TEST_CONDITION( first->GetMITKImage() == reference, "oclImage has the correct reference mitk::Image"); MITK_TEST_CONDITION( first->GetDimension() == reference->GetDimension(), "Same dimensionality."); cl_int clErr = 0; cl_command_queue cmdQueue = clCreateCommandQueue( gpuContext, gpuDevice, 0 ,&clErr); MITK_TEST_CONDITION_REQUIRED( clErr == CL_SUCCESS, "A command queue was created."); // Allocate and copy image data to GPU first->TransferDataToGPU(cmdQueue); MITK_TEST_CONDITION( first->IsModified(0), "Modified flag for GPU correctly set."); // check if the created GPU object is valid cl_mem gpuImage = first->GetGPUImage(cmdQueue); MITK_TEST_CONDITION_REQUIRED( gpuImage != NULL, "oclImage returned a valid GPU memory pointer"); size_t returned = 0; cl_image_format imgFmt; clErr = clGetImageInfo( gpuImage, CL_IMAGE_FORMAT, sizeof(cl_image_format), (void*) &imgFmt, &returned ); MITK_TEST_CONDITION( clErr == CL_SUCCESS, "oclImage has created a valid GPU image"); // test for dimensions size_t imagesize = 0; clErr = clGetImageInfo( gpuImage, CL_IMAGE_WIDTH, sizeof(size_t), (void*) &imagesize, &returned ); MITK_TEST_CONDITION( imagesize == static_cast(first->GetDimension(0)), "Image width corresponds" ); clErr = clGetImageInfo( gpuImage, CL_IMAGE_HEIGHT, sizeof(size_t), (void*) &imagesize, &returned ); MITK_TEST_CONDITION( imagesize == static_cast(first->GetDimension(1)), "Image height corresponds" ); clErr = clGetImageInfo( gpuImage, CL_IMAGE_DEPTH, sizeof(size_t), (void*) &imagesize, &returned ); MITK_TEST_CONDITION( imagesize == static_cast(first->GetDimension(2)), "Image depth corresponds" ); // clean up clReleaseCommandQueue( cmdQueue ); MITK_TEST_END(); -} +} \ No newline at end of file diff --git a/Modules/OpenCL/Testing/mitkOclReferenceCountTest.cpp b/Modules/OpenCL/Testing/mitkOclReferenceCountTest.cpp index bed5d0b950..180c76af11 100644 --- a/Modules/OpenCL/Testing/mitkOclReferenceCountTest.cpp +++ b/Modules/OpenCL/Testing/mitkOclReferenceCountTest.cpp @@ -1,92 +1,92 @@ /*=================================================================== 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 #include #include #include // itk filter for reference computation #include #include #include using namespace mitk; /** This function is testing the OclFilter class and the OpenCL resource service. To prevent segmentation faults a mutexed reference counter is implemented in the resource service. It tracks the number of opencl program references for the corresponding filter and delete only the opencl programm if the reference count reaches 0. Every new instance of a filter increases the reference count by 1. This test runs successfull if the 2 filters are initialized, run and deleted without any crash. */ int mitkOclReferenceCountTest( int argc, char* argv[] ) { MITK_TEST_BEGIN("mitkOclReferenceCountTest"); // instancate uService - ServiceReference ref = GetModuleContext()->GetServiceReference(); + us::ServiceReference ref = GetModuleContext()->GetServiceReference(); OclResourceService* resources = GetModuleContext()->GetService(ref); cl_context gpuContext = resources->GetContext(); cl_device_id gpuDevice = resources->GetCurrentDevice(); //Create a random reference image mitk::Image::Pointer inputImage = mitk::ImageGenerator::GenerateRandomImage(119, 204, 52, 1, // dimension 1.0f, 1.0f, 1.0f, // spacing 255, 0); // max, min int upperThr = 255; int lowerThr = 60; int outsideVal = 0; int insideVal = 100; mitk::OclBinaryThresholdImageFilter* oclFilter1 = new mitk::OclBinaryThresholdImageFilter; oclFilter1->SetInput( inputImage ); oclFilter1->SetUpperThreshold( upperThr ); oclFilter1->SetLowerThreshold( lowerThr ); oclFilter1->SetOutsideValue( outsideVal ); oclFilter1->SetInsideValue( insideVal ); oclFilter1->Update(); mitk::Image::Pointer outputImage1 = mitk::Image::New(); outputImage1 = oclFilter1->GetOutput(); mitk::OclBinaryThresholdImageFilter* oclFilter2 = new mitk::OclBinaryThresholdImageFilter; oclFilter2->SetInput( inputImage ); oclFilter2->SetUpperThreshold( upperThr ); oclFilter2->SetLowerThreshold( lowerThr ); oclFilter2->SetOutsideValue( outsideVal ); oclFilter2->SetInsideValue( insideVal ); oclFilter2->Update(); mitk::Image::Pointer outputImage2 = mitk::Image::New(); outputImage2 = oclFilter2->GetOutput(); // delete filters delete oclFilter1; delete oclFilter2; // this is only visible if the delete did not cause a segmentation fault // it is always true and successfull if the program reaches this state MITK_TEST_CONDITION_REQUIRED( true, "2 Filters deleted without a crash -> success "); MITK_TEST_END(); -} +} \ No newline at end of file diff --git a/Modules/OpenCL/Testing/mitkOclResourceServiceTest.cpp b/Modules/OpenCL/Testing/mitkOclResourceServiceTest.cpp index b20f4bea4f..7ad78da5cf 100644 --- a/Modules/OpenCL/Testing/mitkOclResourceServiceTest.cpp +++ b/Modules/OpenCL/Testing/mitkOclResourceServiceTest.cpp @@ -1,115 +1,115 @@ /*=================================================================== 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 "mitkTestingMacros.h" #include "mitkOclUtils.h" #include #include #include #include "mitkOclResourceService.h" #include "mitkException.h" #include #include using namespace mitk; /** This function is testing the class mitk::OclContextManager. */ int mitkOclResourceServiceTest( int argc, char* argv[] ) { MITK_TEST_BEGIN("mitkOclResourceServiceTest"); - ServiceReference ref = GetModuleContext()->GetServiceReference(); + us::ServiceReference ref = us::GetModuleContext()->GetServiceReference(); MITK_TEST_CONDITION_REQUIRED( ref != NULL, "Resource service available." ); - OclResourceService* resources = GetModuleContext()->GetService(ref); + OclResourceService* resources = us::GetModuleContext()->GetService(ref); MITK_TEST_CONDITION_REQUIRED( resources != NULL, "Resource service available." ); cl_context first = resources->GetContext(); MITK_TEST_CONDITION_REQUIRED(first != NULL, "Got not-null OpenCL context."); - OclResourceService* resources_2 = GetModuleContext()->GetService(ref); + OclResourceService* resources_2 = us::GetModuleContext()->GetService(ref); MITK_TEST_CONDITION_REQUIRED( resources == resources_2, "Same resource reference the second time." ); cl_context second = resources_2->GetContext(); MITK_TEST_CONDITION_REQUIRED( first == second, "Both return same context"); // further tests requires for valid context if( first ) { cl_image_format testFmt; testFmt.image_channel_data_type = CL_FLOAT; testFmt.image_channel_order = CL_RGBA; MITK_TEST_CONDITION( resources->GetIsFormatSupported( &testFmt ), "Checking if format CL_FLOAT / CL_RGBA supported." ); } // create test program const std::string testProgramSource = "__kernel void testKernel( __global uchar* buffer ){ \ const unsigned int globalPosX = get_global_id(0); \ buffer[globalPosX] = buffer[globalPosX] + 1;}"; cl_int err = 0; size_t progSize = testProgramSource.length(); const char* progSource = testProgramSource.c_str(); cl_program testProgram = clCreateProgramWithSource( first, 1, &progSource, &progSize, &err ); MITK_TEST_CONDITION_REQUIRED( err == CL_SUCCESS, "Test program loaded succesfully."); err = clBuildProgram(testProgram, 0, NULL, NULL, NULL, NULL); MITK_TEST_CONDITION_REQUIRED( err == CL_SUCCESS, "Test program built succesfully."); resources->InsertProgram( testProgram, "test_program", true); MITK_TEST_CONDITION( resources->GetProgram("test_program") == testProgram, "Program correctly stored by ResourceService"); // the manger throws exception when accessing non-existant programs MITK_TEST_FOR_EXCEPTION( mitk::Exception, resources->GetProgram("blah"); ); // another test source, this one does not compile const std::string testProgramSource_notCompiling = "__kernel void testKernel( __global uchar* buffer ){ \ const unsigned intt globalPosX = get_global_id(0); }"; progSize = testProgramSource_notCompiling.length(); const char* progSource2 = testProgramSource_notCompiling.c_str(); cl_program notComp_testProgram = clCreateProgramWithSource( first, 1, &progSource2, &progSize, &err ); // the error in the source code has no influence on loading the program MITK_TEST_CONDITION_REQUIRED( err == CL_SUCCESS, "Test program 2 loaded succesfully."); err = clBuildProgram(notComp_testProgram, 0, NULL, NULL, NULL, NULL); MITK_TEST_CONDITION_REQUIRED( err == CL_BUILD_PROGRAM_FAILURE, "Test program 2 failed to build."); std::cout << " --> The (expected) OpenCL Build Error occured : ";// << GetOclErrorString(err); resources->InsertProgram( notComp_testProgram, "test_program_failed", true); MITK_TEST_CONDITION( resources->GetProgram("test_program_failed") == notComp_testProgram, "Program correctly stored by ResourceService"); // calling the InvalidateStorage() will result in removing the _failed test program inserted above resources->InvalidateStorage(); // the second test program should no more exist in the storage, hence we await an exception MITK_TEST_FOR_EXCEPTION( mitk::Exception, resources->GetProgram("test_program_failed"); ); MITK_TEST_END(); } -US_INITIALIZE_MODULE("OpenCLTestDriver", "", "", "" ) +US_INITIALIZE_MODULE("OpenCLTestDriver", "", "", "" ) \ No newline at end of file diff --git a/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.cpp b/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.cpp index 4e7dede7a7..59ec122ff6 100644 --- a/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.cpp +++ b/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.cpp @@ -1,100 +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" +#include "usServiceReference.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(); + us::ServiceReference ref = GetModuleContext()->GetServiceReference(); OclResourceService* resources = GetModuleContext()->GetService(ref); // clean-up also the resources 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 ); -} +} \ No newline at end of file diff --git a/Modules/OpenCL/mitkOclFilter.cpp b/Modules/OpenCL/mitkOclFilter.cpp index 9e61e137da..c8596d73f9 100644 --- a/Modules/OpenCL/mitkOclFilter.cpp +++ b/Modules/OpenCL/mitkOclFilter.cpp @@ -1,224 +1,227 @@ /*=================================================================== 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 +//usService +#include "usServiceReference.h" + 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) { - mitk::ServiceReference ref = GetModuleContext()->GetServiceReference(); + us::ServiceReference ref = GetModuleContext()->GetServiceReference(); OclResourceService* resources = GetModuleContext()->GetService(ref); // remove program from storage resources->RemoveProgram(m_FilterID); } } 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() { - mitk::ServiceReference ref = GetModuleContext()->GetServiceReference(); + us::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 = 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 - mitk::ServiceReference ref = GetModuleContext()->GetServiceReference(); + us::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, resources->GetCurrentDevice() ); oclLogBinary(m_ClProgram, resources->GetCurrentDevice() ); m_Initialized = false; } // store the succesfully build program into the program storage provided by the resource service 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; this->m_GlobalWorkSize[0] = dimx; this->m_GlobalWorkSize[1] = dimy; this->m_GlobalWorkSize[2] = dimz; // estimate the global work size this->m_GlobalWorkSize[0] = iDivUp( dimx, this->m_LocalWorkSize[0]) * this->m_LocalWorkSize[0]; if ( dimy > 1) this->m_GlobalWorkSize[1] = iDivUp( dimy, this->m_LocalWorkSize[1]) * this->m_LocalWorkSize[1]; if( dimz > 1 ) 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; -} +} \ No newline at end of file diff --git a/Modules/OpenCL/mitkOclImage.cpp b/Modules/OpenCL/mitkOclImage.cpp index 7ad3b08c65..230c92334f 100644 --- a/Modules/OpenCL/mitkOclImage.cpp +++ b/Modules/OpenCL/mitkOclImage.cpp @@ -1,355 +1,350 @@ /*=================================================================== 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 #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; iGetServiceReference(); + us::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 ) { mitk::ImageReadAccessor accessor(m_Image); clErr = clEnqueueWriteImage( gpuComQueue, m_gpuImage, CL_TRUE, origin, region, 0, 0, accessor.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; - mitk::ServiceReference ref = GetModuleContext()->GetServiceReference(); + us::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 = 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]; - -} +} \ No newline at end of file diff --git a/Modules/OpenCL/mitkOclResourceServiceImpl_p.h b/Modules/OpenCL/mitkOclResourceServiceImpl_p.h index 72015fcf2f..98b17e0ec1 100644 --- a/Modules/OpenCL/mitkOclResourceServiceImpl_p.h +++ b/Modules/OpenCL/mitkOclResourceServiceImpl_p.h @@ -1,184 +1,186 @@ /*=================================================================== 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 __mitkOclResourceServiceImpl_h #define __mitkOclResourceServiceImpl_h #include //Micro Services #include #include #include #include //ocl #include "mitkOclResourceService.h" #include "mitkOclUtils.h" #include "mitkOclImageFormats.h" #include +US_USE_NAMESPACE + //todo add docu! /** @struct OclContextCollection * @brief An capsulation of all OpenCL context related variables needed for the OclResourceService implementation * * The struct gets created on first call to GetContext in the OclResourceService and attepts to initialize all * relevant parts, i.e. the context itself, the device and the command queue */ struct OclContextCollection{ public: OclContextCollection() : m_Context(NULL), m_Devices(NULL), m_CreateContextFailed(false) { cl_int clErr = 0; size_t szParmDataBytes; cl_platform_id cpPlatform; cl_device_id m_cdDevice; try{ clErr = oclGetPlatformID( &cpPlatform); CHECK_OCL_ERR( clErr ); clErr = clGetDeviceIDs( cpPlatform, CL_DEVICE_TYPE_GPU, 1, &m_cdDevice, NULL); CHECK_OCL_ERR( clErr ); this->m_Context = clCreateContext( 0, 1, &m_cdDevice, NULL, NULL, &clErr); m_CreateContextFailed = (clErr != CL_SUCCESS); // get the info size clErr = clGetContextInfo(m_Context, CL_CONTEXT_DEVICES, 0,NULL, &szParmDataBytes ); this->m_Devices = (cl_device_id*) malloc(szParmDataBytes); // get device info clErr = clGetContextInfo(m_Context, CL_CONTEXT_DEVICES, szParmDataBytes, m_Devices, NULL); CHECK_OCL_ERR( clErr ); // create command queue m_CommandQueue = clCreateCommandQueue(m_Context, m_Devices[0], 0, &clErr); CHECK_OCL_ERR( clErr ); this->PrintContextInfo( ); // collect available image formats for current context this->m_ImageFormats = mitk::OclImageFormats::New(); this->m_ImageFormats->SetGPUContext(m_Context); } catch( std::exception& e) { MITK_ERROR("OpenCL.ResourceService") << "Exception while creating context: \n" << e.what(); } } ~OclContextCollection() { // if devices were allocated, delete if(m_Devices) { // TODO: Available first in OpenCL 1.2 : query the device for CL_PLATFORM_VERSION // through clGetPlatformInfo // clReleaseDevice(m_Devices[0]); delete [] m_Devices; } // if context was created release it if( m_Context ) clReleaseContext( this->m_Context ); } bool CanProvideContext() const { return ( m_Context != NULL && !m_CreateContextFailed ); } void PrintContextInfo() const { if( m_Devices ) { oclPrintDeviceInfo( m_Devices[0] ); } } /** The context */ cl_context m_Context; /** Available OpenCL devices */ cl_device_id* m_Devices; /** Class for handling (un)supported GPU image formats **/ mitk::OclImageFormats::Pointer m_ImageFormats; /** The command queue*/ cl_command_queue m_CommandQueue; bool m_CreateContextFailed; }; class OclResourceServiceImpl - : public US_BASECLASS_NAME, public OclResourceService + : public OclResourceService { private: // define programmdata private class struct ProgramData { int counter; cl_program program; itk::FastMutexLock::Pointer mutex; ProgramData() :counter(1), program(NULL) {} }; typedef std::map< std::string, ProgramData > ProgramMapType; //typedef std::map< std::string, std::pair< int, cl_program> > ProgramMapType; mutable OclContextCollection* m_ContextCollection; /** Map containing all available (allready compiled) OpenCL Programs */ ProgramMapType m_ProgramStorage; /** mutex for manipulating the program storage */ itk::FastMutexLock::Pointer m_ProgramStorageMutex; public: OclResourceServiceImpl(); ~OclResourceServiceImpl(); cl_context GetContext() const; cl_command_queue GetCommandQueue() const; cl_device_id GetCurrentDevice() const; bool GetIsFormatSupported(cl_image_format *); void PrintContextInfo() const; void InsertProgram(cl_program _program_in, std::string name, bool forceOverride=true); cl_program GetProgram(const std::string&name); void InvalidateStorage(); void RemoveProgram(const std::string&name); unsigned int GetMaximumImageSize(unsigned int dimension, cl_mem_object_type _imagetype); }; #endif // __mitkOclResourceServiceImpl_h diff --git a/Modules/OpenCL/mitkOpenCLActivator.cpp b/Modules/OpenCL/mitkOpenCLActivator.cpp index 29e2efb2c0..5018bb678c 100644 --- a/Modules/OpenCL/mitkOpenCLActivator.cpp +++ b/Modules/OpenCL/mitkOpenCLActivator.cpp @@ -1,33 +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) +void OpenCLActivator::Load(us::ModuleContext *context) { // generate context m_ResourceService.reset(new OclResourceServiceImpl); - mitk::ServiceProperties props; + us::ServiceProperties props; context->RegisterService(m_ResourceService.get(), props); } -void OpenCLActivator::Unload(mitk::ModuleContext *) +void OpenCLActivator::Unload(us::ModuleContext *) { m_ResourceService.release(); } -US_EXPORT_MODULE_ACTIVATOR(MitkOcl, OpenCLActivator ) +US_EXPORT_MODULE_ACTIVATOR(MitkOcl, OpenCLActivator ) \ No newline at end of file diff --git a/Modules/OpenCL/mitkOpenCLActivator.h b/Modules/OpenCL/mitkOpenCLActivator.h index c3132404d0..35308ce002 100644 --- a/Modules/OpenCL/mitkOpenCLActivator.h +++ b/Modules/OpenCL/mitkOpenCLActivator.h @@ -1,54 +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 +class US_ABI_LOCAL OpenCLActivator : public us::ModuleActivator { private: std::auto_ptr m_ResourceService; public: /** @brief Load module context */ - void Load(mitk::ModuleContext *context); + void Load(us::ModuleContext *context); /** @brief Unload module context */ - void Unload(mitk::ModuleContext* ); + void Unload(us::ModuleContext* ); }; #endif // __mitkOpenCLActivator_h