diff --git a/Modules/OpenCL/Resources/BinaryThresholdFilter.cl b/Modules/OpenCL/Resources/BinaryThresholdFilter.cl new file mode 100644 index 0000000000..1c0056eb3f --- /dev/null +++ b/Modules/OpenCL/Resources/BinaryThresholdFilter.cl @@ -0,0 +1,53 @@ +/*=================================================================== + +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. + +===================================================================*/ + +__kernel void ckBinaryThreshold( + __read_only image3d_t dSource, // input image + __global uchar* dDest, // output buffer + int lowerT, int upperT, int outsideVal, int insideVal // parameters +) +{ + // get thread identifier + unsigned int globalPosX = get_global_id(0); + unsigned int globalPosY = get_global_id(1); + unsigned int globalPosZ = get_global_id(2); + + // get image width and weight + const unsigned int uiWidth = get_image_width( dSource ); + const unsigned int uiHeight = get_image_height( dSource ); + const unsigned int uiDepth = get_image_depth( dSource ); + + // create an image sampler + const sampler_t defaultSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST ; + + // terminate non-valid threads + if ( globalPosX < uiWidth && globalPosY < uiHeight && globalPosZ < uiDepth ) + { + int result = outsideVal; + + // get input value + int4 pixelValue = read_imagei( dSource, defaultSampler, (int4) (globalPosX, globalPosY, globalPosZ, 0 )); + + if( (pixelValue.x >= lowerT) && (pixelValue.x <= upperT ) ) + { + result = insideVal; + } + + // store the result + dDest[ globalPosZ * uiWidth * uiHeight + globalPosY * uiWidth + globalPosX ] = result; + } + +} diff --git a/Modules/OpenCL/files.cmake b/Modules/OpenCL/files.cmake index aa0d3ecca4..799077854e 100644 --- a/Modules/OpenCL/files.cmake +++ b/Modules/OpenCL/files.cmake @@ -1,19 +1,23 @@ set(CPP_FILES # helper classes mitkOclUtils.cpp mitkOclResourceServiceImpl_Private.cpp mitkOclImageFormats.cpp # module activator mitkOpenCLActivator.cpp # base data and filter objects mitkOclBaseData.cpp mitkOclImage.cpp mitkOclFilter.cpp mitkOclImageFilter.cpp mitkOclImageToImageFilter.cpp # own filter implementations mitkOclBinaryThresholdImageFilter.cpp ) + +set(RESOURCE_FILES + BinaryThresholdFilter.cl +) diff --git a/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.cpp b/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.cpp index 59ec122ff6..c451265b37 100644 --- a/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.cpp +++ b/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.cpp @@ -1,100 +1,104 @@ /*=================================================================== 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->AddSourceFile("BinaryThresholdFilter.cl"); 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()) { 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 ); } +us::Module *mitk::OclBinaryThresholdImageFilter::GetModule() +{ + return us::GetModuleContext()->GetModule(); +} + 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/mitkOclBinaryThresholdImageFilter.h b/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.h index a780d3c933..dc3168c284 100644 --- a/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.h +++ b/Modules/OpenCL/mitkOclBinaryThresholdImageFilter.h @@ -1,112 +1,114 @@ /*=================================================================== 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 _MITKOCLBINARYTHRESHOLDIMAGEFILTER_H_ #define _MITKOCLBINARYTHRESHOLDIMAGEFILTER_H_ #include "mitkOclImageToImageFilter.h" namespace mitk { class OclImageToImageFilter; /** Documentation * * \brief The OclBinaryThresholdImageFilter computes a binary segmentation based on given threshold values. * * The filter requires two threshold values ( the upper and the lower threshold ) and two image values ( inside and outside ). The resulting voxel of the segmentation image is assigned the inside value 1 if the image value is between the given thresholds and the outside value otherwise. */ class MitkOcl_EXPORT OclBinaryThresholdImageFilter : public OclImageToImageFilter { typedef OclFilter Superclass; public: /** Update the filter */ void Update(); /** Constructor */ OclBinaryThresholdImageFilter(); /** Destructor */ virtual ~OclBinaryThresholdImageFilter(); /** Set the lower threshold @param thr Threshold value */ void SetLowerThreshold( int lowerThreshold ) { this->m_LowerThreshold = lowerThreshold; } /** Set the upper threshold @param thr Threshold value */ void SetUpperThreshold( int upperThreshold ) { this->m_UpperThreshold = upperThreshold; } /** Set the outside value @param val The outside value */ void SetOutsideValue( int outsideValue ) { this->m_OutsideValue = outsideValue; } /** Set the inside value @param val The inside value */ void SetInsideValue( int insideValue ) { this->m_InsideValue = insideValue; } protected: /** Initialize the filter */ bool Initialize(); void Execute(); mitk::PixelType GetOutputType() { return mitk::MakeScalarPixelType(); } int GetBytesPerElem() { return sizeof(unsigned char); } + virtual us::Module* GetModule(); + private: /** The OpenCL kernel for the filter */ cl_kernel m_ckBinaryThreshold; int m_LowerThreshold; int m_UpperThreshold; int m_InsideValue; int m_OutsideValue; }; } #endif diff --git a/Modules/OpenCL/mitkOclFilter.cpp b/Modules/OpenCL/mitkOclFilter.cpp index c8596d73f9..a726d63b86 100644 --- a/Modules/OpenCL/mitkOclFilter.cpp +++ b/Modules/OpenCL/mitkOclFilter.cpp @@ -1,227 +1,237 @@ /*=================================================================== 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" +#include +#include +#include +#include +#include +#include mitk::OclFilter::OclFilter() - : m_ClFile(), - m_ClSource(NULL), - m_ClCompilerFlags(""), + : 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_ClCompilerFlags(""), m_ClProgram(NULL), m_CommandQue(NULL), m_FilterID(filename), m_Preambel(" "), m_Initialized(false) { - m_ClSourcePath = MITK_ROOT; - m_ClSourcePath += "Modules/OpenCL/ShaderSources"; + m_ClFiles.push_back(filename); } mitk::OclFilter::~OclFilter() { MITK_DEBUG << "OclFilter Destructor"; // release program if (m_ClProgram) { 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() { 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())) + if ( m_ClFiles.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) +void mitk::OclFilter::LoadSourceFiles(CStringList &sourceCode, ClSizeList &sourceCodeSize) { - MITK_DEBUG("ocl.filter") << "Setting source [" << filename <<" ]"; + for( CStringList::iterator it = m_ClFiles.begin(); it != m_ClFiles.end(); ++it ) + { + + MITK_INFO << "File :" << *it; + us::ModuleResource mdr = GetModule()->GetResource(*it); + + if( !mdr.IsValid() ) + MITK_WARN << "Could not load resource: " << mdr.GetName() << " is invalid!"; + + us:ModuleResourceStream rss(mdr); + + std::istreambuf_iterator eos; + std::string src(std::istreambuf_iterator(rss), eos); + + char* tmp = new char[src.size() + 1]; - mitk::StandardFileLocations::GetInstance()->AddDirectoryForSearch( m_ClSourcePath.c_str(), true); - // search for file - m_ClFile = mitk::StandardFileLocations::GetInstance()->FindFile( filename); + strcpy(tmp,src.c_str()); + + sourceCode.push_back(tmp); + sourceCodeSize.push_back(src.size()); + } } void mitk::OclFilter::CompileSource() { - if (m_ClFile.empty() && m_ClSource == NULL) + if (m_ClFiles.empty()) { MITK_ERROR("ocl.filter") << "No shader source file was set"; return; } - // help variables - size_t szKernelLength; + // helper variable + CStringList sourceCode; + ClSizeList sourceCodeSize; int clErr = 0; //get a valid opencl context 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); + LoadSourceFiles(sourceCode, sourceCodeSize); - if (m_ClSource != NULL) + if ( !sourceCode.empty() ) { - m_ClProgram = clCreateProgramWithSource(gpuContext, 1, (const char**)&m_ClSource, &szKernelLength, &clErr); + m_ClProgram = clCreateProgramWithSource(gpuContext, sourceCode.size(), &sourceCode[0], &sourceCodeSize[0], &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) +void mitk::OclFilter::AddSourceFile(const char* filename) { - m_ClSourcePath = path; + m_ClFiles.push_back(filename); } 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/mitkOclFilter.h b/Modules/OpenCL/mitkOclFilter.h index c8094a3739..1c8c7c1029 100644 --- a/Modules/OpenCL/mitkOclFilter.h +++ b/Modules/OpenCL/mitkOclFilter.h @@ -1,146 +1,139 @@ /*=================================================================== 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 __mitkOclFilter_h #define __mitkOclFilter_h #include "mitkOclUtils.h" #include "mitkCommon.h" +#include +#include + #include namespace mitk { /** @class OclFilter @brief Superclass for all OpenCL based filter. This class takes care of loading and compiling the external GPU program code. */ class MitkOcl_EXPORT OclFilter { public: /** * @brief Set the source file of the OpenCL shader * * @param filename Path to the file */ - void SetSourceFile(const char* filename); - - /** - * @brief Set the folder of the directory that contains - * the OpenCL source files. This location will also be used - * as an include folder for the OpenCL compiler and all headers placed - * there can be loaded by the compiler. - * - * @param path to the modulefolder that contains the gpuSource - */ - void SetSourcePath(const char* path); +// void SetSourceFile(const char* filename); + void AddSourceFile(const char* filename); /** * @brief Set specific compilerflags to compile the CL source. Default is set to NULL; * example: "-cl-fast-relaxed-math -cl-mad-enable -cl-strict-aliasing" * * @param flags to the modulefolder that contains the gpuSource */ void SetCompilerFlags(const char* flags); /** @brief Returns true if the initialization was successfull */ virtual bool IsInitialized(); /** @brief Destructor */ virtual ~OclFilter(); protected: + + typedef std::vector CStringList; + typedef std::vector ClSizeList; + /** @brief Constructor */ OclFilter(); /** @brief Constructor ( overloaded ) */ OclFilter(const char* filename); - /** @brief Path to the *.cl source file */ - std::string m_ClFile; - - /** @brief The source code to be compiled for the GPU */ - const char* m_ClSource; - /** @brief String that contains the compiler flags */ const char* m_ClCompilerFlags; /** @brief The compiled OpenCL program */ cl_program m_ClProgram; /** @brief Command queue for the filter */ cl_command_queue m_CommandQue; /** @brief Unique ID of the filter, needs to be specified in the constructor of the derived class */ std::string m_FilterID; /*! @brief source preambel for e.g. #define commands to be inserted into the OpenCL source */ const char* m_Preambel; + CStringList m_ClFiles; + /** @brief status of the filter */ bool m_Initialized; - /** @brief The path of the module folder in wich the gpu sourcefiles are stored */ - std::string m_ClSourcePath; - /** @brief The local work size fo the filter */ size_t m_LocalWorkSize[3]; /** @brief The global work size of the filter */ size_t m_GlobalWorkSize[3]; /** @brief Set the working size for the following OpenCL kernel call */ void SetWorkingSize(unsigned int locx, unsigned int dimx, unsigned int locy = 1, unsigned int dimy = 1, unsigned int locz = 1, unsigned int dimz = 1); /** @brief Execute the given kernel on the OpenCL Index-Space defined by the local and global work sizes */ bool ExecuteKernel( cl_kernel kernel, unsigned int workSizeDim ); /** * \brief Initialize all necessary parts of the filter * * The Initialize() method creates the command queue and the m_clProgram. * The program is either compiled from the given source or taken from the * OclResourceManager if the program was compiled already. */ bool Initialize(); /** * @brief Compile the program source * * @param preambel e.g. defines for the shader code */ void CompileSource(); /** * @brief Add some source code on the beginning of the loaded source * * In this way, some preprocessor flags for the CL compiler can at the beginning of the filter * @param preambel Source preambel for e.g. #define commands to be inserted into the OpenCL source */ void SetSourcePreambel(const char* preambel); + virtual us::Module* GetModule() = 0; + void LoadSourceFiles(CStringList &SourceCodeList, ClSizeList &SourceCodeSizeList); }; } #endif // __mitkOclFilter_h diff --git a/Modules/OpenCL/mitkOclUtils.cpp b/Modules/OpenCL/mitkOclUtils.cpp index 7459941f07..5b572d5b64 100644 --- a/Modules/OpenCL/mitkOclUtils.cpp +++ b/Modules/OpenCL/mitkOclUtils.cpp @@ -1,620 +1,569 @@ /*=================================================================== 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 "mitkOclUtils.h" #include "mitkLogMacros.h" #include #include unsigned int iDivUp(unsigned int dividend, unsigned int divisor){ return (dividend % divisor == 0) ? (dividend / divisor) : (dividend / divisor + 1); } cl_int oclGetPlatformID(cl_platform_id* selectedPlatform) { cl_uint num_platforms = 0; cl_platform_id* clPlatformIDs; cl_int ciErrNum = 0; ciErrNum = clGetPlatformIDs( 0, NULL, &num_platforms); if ( ciErrNum != CL_SUCCESS) { MITK_ERROR<<" Error " << ciErrNum << " in clGetPlatformIDs() \n"; throw std::bad_exception(); } else { clPlatformIDs = new cl_platform_id[num_platforms]; ciErrNum = clGetPlatformIDs( num_platforms, clPlatformIDs, NULL); if(ciErrNum == CL_SUCCESS) { *selectedPlatform = clPlatformIDs[0]; } } return CL_SUCCESS; } void oclPrintMemObjectInfo(cl_mem memobj) { cl_int clErr = 0; MITK_INFO << "Examining cl_mem object: " << memobj << "\n------------------\n"; // CL_MEM_TYPE cl_mem_object_type objtype; clErr = clGetMemObjectInfo( memobj, CL_MEM_TYPE, sizeof(cl_mem_object_type),&objtype, NULL); CHECK_OCL_ERR( clErr ); switch(objtype) { case CL_MEM_OBJECT_BUFFER: MITK_INFO << "CL_MEM_TYPE \t" << "BUFFER_OBJ" << "\n"; break; case CL_MEM_OBJECT_IMAGE2D: MITK_INFO << "CL_MEM_TYPE \t" << "2D IMAGE" << "\n"; break; case CL_MEM_OBJECT_IMAGE3D: MITK_INFO << "CL_MEM_TYPE \t" << "3D IMAGE" << "\n"; break; default: MITK_INFO << "CL_MEM_TYPE \t" << "[could not resolve]" << "\n"; break; } // CL_MEM_FLAGS cl_mem_flags flags; clErr = clGetMemObjectInfo( memobj, CL_MEM_FLAGS, sizeof(cl_mem_flags),&flags, NULL); CHECK_OCL_ERR( clErr ); switch(flags) { case CL_MEM_READ_ONLY: MITK_INFO << "CL_MEM_FLAGS \t" << "CL_MEM_READ_ONLY" << "\n"; break; case CL_MEM_WRITE_ONLY: MITK_INFO << "CL_MEM_FLAGS \t" << "CL_MEM_WRITE_ONLY" << "\n"; break; case CL_MEM_READ_WRITE: MITK_INFO << "CL_MEM_FLAGS \t" << "CL_MEM_READ_WRITE" << "\n"; break; default: MITK_INFO << "CL_MEM_FLAGS \t" << "not resolved, " << flags << "\n"; break; } // get CL_MEM_SIZE size_t memsize; clErr = clGetMemObjectInfo( memobj, CL_MEM_SIZE, sizeof(memsize),&memsize, NULL); CHECK_OCL_ERR( clErr ); MITK_INFO << "CL_MEM_SIZE \t" << memsize << "\n"; // get CL_MEM_HOST_PTR float *hostptr; clErr = clGetMemObjectInfo( memobj, CL_MEM_HOST_PTR, sizeof(void*), (void*) &hostptr, NULL); CHECK_OCL_ERR( clErr ); MITK_INFO << "CL_MEM_HOST_PTR \t" << hostptr << "\n"; // get CL_CONTEXT cl_context gpuctxt; clErr = clGetMemObjectInfo( memobj, CL_MEM_CONTEXT, sizeof(cl_context), &gpuctxt, NULL); CHECK_OCL_ERR( clErr ); MITK_INFO << "CL_CONTEXT \t\t" << gpuctxt << "\n"; // get CL_MEM_REFERENCE_COUNT cl_uint refs; clErr = clGetMemObjectInfo( memobj, CL_MEM_REFERENCE_COUNT, sizeof(cl_uint), &refs, NULL); CHECK_OCL_ERR(clErr); MITK_INFO << "CL_REF_COUNT \t" << refs << "\n"; MITK_INFO << "================== \n" << std::endl; } - -char* oclLoadProgramSource( const char* srcFilename, const char* srcPreamble, size_t* srcLength) -{ - FILE* pFileStream = NULL; - size_t szSourceLength; - - MITK_INFO<< "Loading GPU SourceCode from " << srcFilename <<" \n"; - -#ifdef _WIN32 - if(fopen_s(&pFileStream, srcFilename, "rb") != 0) - { - return NULL; - printf("Failed to open file \n"); - - } -#else // Linux version - pFileStream = fopen(srcFilename, "rb"); - if(pFileStream == 0) - { - return NULL; - } -#endif - - size_t szPreambleLength = strlen(srcPreamble); - - // get the length of the source code - fseek(pFileStream, 0, SEEK_END); - szSourceLength = ftell(pFileStream); - fseek(pFileStream, 0, SEEK_SET); - - // allocate a buffer for the source code string and read it in - char* cSourceString = (char *)malloc(szSourceLength + szPreambleLength + 1); - memcpy(cSourceString, srcPreamble, szPreambleLength); - if (fread((cSourceString) + szPreambleLength, szSourceLength, 1, pFileStream) != 1) - { - fclose(pFileStream); - free(cSourceString); - return 0; - } - - // close the file and return the total length of the combined (preamble + source) string - fclose(pFileStream); - if(srcLength != 0) - { - *srcLength = szSourceLength + szPreambleLength; - } - cSourceString[szSourceLength + szPreambleLength] = '\0'; - - return cSourceString; -} - void oclPrintDeviceInfo(cl_device_id device) { char device_string[1024]; clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(device_string), &device_string, NULL); MITK_INFO("ocl.log")<< " Device : " << device_string; // CL_DEVICE_INFO cl_device_type type; clGetDeviceInfo(device, CL_DEVICE_TYPE, sizeof(type), &type, NULL); if( type & CL_DEVICE_TYPE_CPU ) MITK_INFO("ocl.log")<<" CL_DEVICE_TYPE: CL_DEVICE_TYPE_CPU"; if( type & CL_DEVICE_TYPE_GPU ) MITK_INFO("ocl.log")<<" CL_DEVICE_TYPE: CL_DEVICE_TYPE_GPU"; if( type & CL_DEVICE_TYPE_ACCELERATOR ) MITK_INFO("ocl.log")<<" CL_DEVICE_TYPE: CL_DEVICE_TYPE_ACCELERATOR"; if( type & CL_DEVICE_TYPE_DEFAULT ) MITK_INFO("ocl.log")<<" CL_DEVICE_TYPE: CL_DEVICE_TYPE_DEFAULT"; // CL_DEVICE_MAX_COMPUTE_UNITS cl_uint compute_units; clGetDeviceInfo(device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(compute_units), &compute_units, NULL); MITK_INFO("ocl.log")<<" CL_DEVICE_MAX_COMPUTE_UNITS:" << compute_units; // CL_DEVICE_MAX_WORK_GROUP_SIZE size_t workitem_size[3]; clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(workitem_size), &workitem_size, NULL); MITK_INFO("ocl.log")<<" CL_DEVICE_MAX_WORK_ITEM_SIZES:\t"<< workitem_size[0]<< workitem_size[1]<< workitem_size[2]; // CL_DEVICE_MAX_WORK_GROUP_SIZE size_t workgroup_size; clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(workgroup_size), &workgroup_size, NULL); MITK_INFO("ocl.log")<<" CL_DEVICE_MAX_WORK_GROUP_SIZE:" << workgroup_size; // CL_DEVICE_MAX_CLOCK_FREQUENCY cl_uint clock_frequency; clGetDeviceInfo(device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(clock_frequency), &clock_frequency, NULL); MITK_INFO("ocl.log")<<" CL_DEVICE_MAX_CLOCK_FREQUENCY:"<< clock_frequency / 1000; // CL_DEVICE_IMAGE_SUPPORT cl_bool image_support; clGetDeviceInfo(device, CL_DEVICE_IMAGE_SUPPORT, sizeof(image_support), &image_support, NULL); MITK_INFO("ocl.log")<<" CL_DEVICE_IMAGE_SUPPORT:\t" << image_support; // CL_DEVICE_GLOBAL_MEM_SIZE cl_ulong mem_size; clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(mem_size), &mem_size, NULL); MITK_INFO("ocl.log")<<" CL_DEVICE_GLOBAL_MEM_SIZE:\t\t"<<(unsigned int)(mem_size / (1024 * 1024))<<"Mbytes"; // CL_DEVICE_LOCAL_MEM_SIZE clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(mem_size), &mem_size, NULL); MITK_INFO("ocl.log")<<" CL_DEVICE_LOCAL_MEM_SIZE:\t\t"<< (unsigned int)(mem_size / (1024)) <<"KByte\n"; //check for image support properties clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(workgroup_size), &workgroup_size, NULL); MITK_INFO("ocl.log")<<" CL_DEVICE_IMAGE2D_MAX_WIDTH:\t" << workgroup_size; clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(workgroup_size), &workgroup_size, NULL); MITK_INFO("ocl.log")<<" CL_DEVICE_IMAGE2D_MAX_HEIGHT:\t" << workgroup_size; clGetDeviceInfo(device, CL_DEVICE_IMAGE3D_MAX_WIDTH, sizeof(workgroup_size), &workgroup_size, NULL); MITK_INFO("ocl.log")<<" CL_DEVICE_IMAGE3D_MAX_WIDTH:\t" << workgroup_size; clGetDeviceInfo(device, CL_DEVICE_IMAGE3D_MAX_HEIGHT, sizeof(workgroup_size), &workgroup_size, NULL); MITK_INFO("ocl.log")<<" CL_DEVICE_IMAGE3D_MAX_HEIGHT:\t" << workgroup_size; clGetDeviceInfo(device, CL_DEVICE_IMAGE3D_MAX_DEPTH, sizeof(workgroup_size), &workgroup_size, NULL); MITK_INFO("ocl.log")<<" CL_DEVICE_IMAGE3D_MAX_DEPTH:\t" << workgroup_size; // CL_DEVICE_QUEUE_PROPERTIES cl_command_queue_properties queue_properties; clGetDeviceInfo(device, CL_DEVICE_QUEUE_PROPERTIES, sizeof(queue_properties), &queue_properties, NULL); if( queue_properties & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE ) MITK_INFO("ocl.log")<<" CL_DEVICE_QUEUE_PROPERTIES:\t\t"<< "CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE"; if( queue_properties & CL_QUEUE_PROFILING_ENABLE ) MITK_INFO("ocl.log")<<" CL_DEVICE_QUEUE_PROPERTIES:\t\t"<< "CL_QUEUE_PROFILING_ENABLE"; } std::string GetOclErrorAsString( int _clErr ) { std::string returnString("CL_SUCCESS\n"); switch(_clErr) { case CL_DEVICE_NOT_FOUND: returnString = "CL_DEVICE_NOT_FOUND\n"; break; case CL_DEVICE_NOT_AVAILABLE: returnString = "CL_DEVICE_NOT_AVAILABLE\n"; break; /*case CL_DEVICE_COMPILER_NOT_AVAILABLE: returnString = "CL_DEVICE_COMPILER_NOT_AVAILABLE\n"; break; */ case CL_MEM_OBJECT_ALLOCATION_FAILURE : returnString = "CL_MEM_OBJECT_ALLOCATION_FAILURE\n"; break; case CL_OUT_OF_RESOURCES: returnString = "CL_OUT_OF_RESOURCES\n"; break; case CL_OUT_OF_HOST_MEMORY: returnString = "CL_OUT_OF_HOST_MEMORY\n"; break; case CL_PROFILING_INFO_NOT_AVAILABLE: returnString = "CL_PROFILING_INFO_NOT_AVAILABLE\n"; break; case CL_MEM_COPY_OVERLAP: returnString = "CL_MEM_COPY_OVERLAP\n"; break; case CL_IMAGE_FORMAT_MISMATCH: returnString = "CL_IMAGE_FORMAT_MISMATCH\n"; break; case CL_IMAGE_FORMAT_NOT_SUPPORTED: returnString = "CL_IMAGE_FORMAT_NOT_SUPPORTED\n"; break; case CL_BUILD_PROGRAM_FAILURE: returnString = "CL_BUILD_PROGRAM_FAILURE\n"; break; case CL_MAP_FAILURE: returnString = "CL_MAP_FAILURE\n"; break; case CL_INVALID_VALUE: returnString = "CL_INVALID_VALUE\n"; break; case CL_INVALID_DEVICE_TYPE: returnString = "CL_INVALID_DEVICE_TYPE\n"; break; case CL_INVALID_PLATFORM: returnString = "CL_INVALID_PLATFORM\n"; break; case CL_INVALID_DEVICE: returnString = "CL_INVALID_DEVICE\n"; break; case CL_INVALID_CONTEXT : returnString = "CL_INVALID_CONTEXT\n"; break; case CL_INVALID_QUEUE_PROPERTIES: returnString = "CL_INVALID_QUEUE_PROPERTIES\n"; break; case CL_INVALID_COMMAND_QUEUE: returnString = "CL_INVALID_COMMAND_QUEUE\n"; break; case CL_INVALID_HOST_PTR: returnString = "CL_INVALID_HOST_PTR\n"; break; case CL_INVALID_MEM_OBJECT: returnString = "CL_INVALID_MEM_OBJECT\n"; break; case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: returnString = "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR\n"; break; case CL_INVALID_IMAGE_SIZE: returnString = "CL_INVALID_IMAGE_SIZE\n"; break; case CL_INVALID_SAMPLER : returnString = "CL_INVALID_SAMPLER\n"; break; case CL_INVALID_BINARY: returnString = "CL_INVALID_BINARY\n"; break; case CL_INVALID_BUILD_OPTIONS: returnString = "CL_INVALID_BUILD_OPTIONS\n"; break; case CL_INVALID_PROGRAM: returnString = "CL_INVALID_PROGRAM\n"; break; case CL_INVALID_PROGRAM_EXECUTABLE: returnString = "CL_INVALID_PROGRAM_EXECUTABLE\n"; break; case CL_INVALID_KERNEL_NAME: returnString = "CL_INVALID_KERNEL_NAME\n"; break; case CL_INVALID_KERNEL_DEFINITION: returnString = "CL_INVALID_KERNEL_DEFINITION\n"; break; case CL_INVALID_KERNEL : returnString = "CL_INVALID_KERNEL\n"; break; case CL_INVALID_ARG_INDEX : returnString = "CL_INVALID_ARG_INDEX\n"; break; case CL_INVALID_ARG_VALUE : returnString = "CL_INVALID_ARG_VALUE\n"; break; case CL_INVALID_ARG_SIZE : returnString = "CL_INVALID_ARG_SIZE\n"; break; case CL_INVALID_KERNEL_ARGS : returnString = "CL_INVALID_KERNEL_ARGS\n"; break; case CL_INVALID_WORK_DIMENSION: returnString = "CL_INVALID_WORK_DIMENSION\n"; break; case CL_INVALID_WORK_GROUP_SIZE: returnString = "CL_INVALID_WORK_GROUP_SIZE\n"; break; case CL_INVALID_WORK_ITEM_SIZE: returnString = "CL_INVALID_WORK_ITEM_SIZE\n"; break; case CL_INVALID_GLOBAL_OFFSET: returnString = "CL_INVALID_GLOBAL_OFFSET\n"; break; case CL_INVALID_EVENT_WAIT_LIST: returnString = "CL_INVALID_EVENT_WAIT_LIST\n"; break; case CL_INVALID_EVENT: returnString = "CL_INVALID_EVENT\n"; break; case CL_INVALID_OPERATION: returnString = "CL_INVALID_OPERATION\n"; break; case CL_INVALID_GL_OBJECT: returnString = "CL_INVALID_GL_OBJECT\n"; break; case CL_INVALID_BUFFER_SIZE : returnString = "CL_INVALID_BUFFER_SIZE\n"; break; case CL_INVALID_MIP_LEVEL : returnString = "CL_INVALID_MIP_LEVEL\n"; break; default: break; } return returnString; } void GetOclError(int _clErr) { if(_clErr == CL_SUCCESS ) MITK_WARN << "Called GetOclErr() with no error value: [CL_SUCCESS]"; else MITK_ERROR << GetOclErrorAsString(_clErr); } bool oclCheckError(int _err, const char* filepath, int lineno) { if (_err) { MITK_ERROR<< "OpenCL Error at " << filepath <<":"<< lineno; GetOclError(_err); return 0; } return 1; } void GetSupportedImageFormats(cl_context _context, cl_mem_object_type _type) { const unsigned int entries = 500; cl_image_format* formats = new cl_image_format[entries]; cl_uint _written = 0; // OpenCL constant to catch error IDs cl_int ciErr1; // Get list of supported image formats for READ_ONLY access ciErr1 = clGetSupportedImageFormats( _context, CL_MEM_READ_ONLY, _type, entries, formats, &_written); CHECK_OCL_ERR(ciErr1); MITK_INFO << "Supported Image Formats, Image: CL_MEM_READ_ONLY \n"; for (unsigned int i=0; i<_written; i++) { MITK_INFO<< "ChannelType: " << GetImageTypeAsString(formats[i].image_channel_data_type) << "| ChannelOrder: "<< GetImageTypeAsString(formats[i].image_channel_order) <<"\n"; } _written = 0; // Get list of supported image formats for READ_WRITE access ciErr1 = clGetSupportedImageFormats( _context, CL_MEM_READ_WRITE, _type, entries, formats, &_written); CHECK_OCL_ERR(ciErr1); MITK_INFO << "Supported Image Formats, Image: CL_MEM_READ_WRITE (found: " << _written <<") \n"; for (unsigned int i=0; i<_written; i++) { MITK_INFO<< "ChannelType: " << GetImageTypeAsString(formats[i].image_channel_data_type) << "| ChannelOrder: "<< GetImageTypeAsString(formats[i].image_channel_order) <<"\n"; } _written = 0; // Get list of supported image formats for WRITE_ONLY access ciErr1 = clGetSupportedImageFormats( _context, CL_MEM_WRITE_ONLY, _type, entries, formats, &_written); CHECK_OCL_ERR(ciErr1); MITK_INFO << "Supported Image Formats, Image: CL_MEM_WRITE_ONLY (found: " << _written <<") \n"; for (unsigned int i=0; i<_written; i++) { MITK_INFO<< "ChannelType: " << GetImageTypeAsString(formats[i].image_channel_data_type) << "| ChannelOrder: "<< GetImageTypeAsString(formats[i].image_channel_order) <<"\n"; } } std::string GetImageTypeAsString( const unsigned int _in) { switch(_in) { case CL_R: return "CL_R "; break; case CL_A: return "CL_A "; break; case CL_RG: return "CL_RG "; break; case CL_RA: return "CL_RA "; break; case CL_RGB: return "CL_RGB "; break; case CL_RGBA: return "CL_RGBA "; break; case CL_BGRA: return "CL_BGRA "; break; case CL_ARGB: return "CL_ARGB "; break; case CL_INTENSITY: return "CL_INTENSITY "; break; case CL_LUMINANCE: return "CL_LUMINANCE "; break; case CL_SNORM_INT8: return "CL_SNORM_INT8 "; break; case CL_SNORM_INT16: return "CL_SNORM_INT16 "; break; case CL_UNORM_INT8: return "CL_UNORM_INT8 "; break; case CL_UNORM_INT16: return "CL_UNORM_INT16 "; break; case CL_UNORM_SHORT_565: return "CL_UNORM_SHORT_565 "; break; case CL_UNORM_SHORT_555: return "CL_UNORM_SHORT_555 "; break; case CL_UNORM_INT_101010: return "CL_UNORM_INT_101010 "; break; case CL_SIGNED_INT8: return "CL_SIGNED_INT8 "; break; case CL_SIGNED_INT16: return "CL_SIGNED_INT16 "; break; case CL_SIGNED_INT32: return "CL_SIGNED_INT32 "; break; case CL_UNSIGNED_INT8: return "CL_UNSIGNED_INT8 "; break; case CL_UNSIGNED_INT16: return "CL_UNSIGNED_INT16 "; break; case CL_UNSIGNED_INT32: return "CL_UNSIGNED_INT32 "; break; case CL_HALF_FLOAT: return "CL_HALF_FLOAT "; break; case CL_FLOAT: return "CL_FLOAT "; break; default: return "--"; break; } } void oclLogBinary(cl_program clProg, cl_device_id clDev) { // Grab the number of devices associated with the program cl_uint num_devices; clGetProgramInfo(clProg, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &num_devices, NULL); // Grab the device ids cl_device_id* devices = (cl_device_id*) malloc(num_devices * sizeof(cl_device_id)); clGetProgramInfo(clProg, CL_PROGRAM_DEVICES, num_devices * sizeof(cl_device_id), devices, 0); // Grab the sizes of the binaries size_t* binary_sizes = (size_t*)malloc(num_devices * sizeof(size_t)); clGetProgramInfo(clProg, CL_PROGRAM_BINARY_SIZES, num_devices * sizeof(size_t), binary_sizes, NULL); // Now get the binaries char** ptx_code = (char**)malloc(num_devices * sizeof(char*)); for( unsigned int i=0; i #include #define CHECK_OCL_ERR(_er) oclCheckError(_er, __FILE__, __LINE__); /** @brief Method to estimate an integer quotient C from given dividend and divisor higher or equal to the corresponding floating quotient If the divisor is a factor of the dividend, the dividend/divisor is an integer value and is returned. If not, the nearest higher integer is returned. So it holds for the return value C that C * divisor is equal or greater then the dividend. In OpenCL context useful for estimating the local/global working dimension of a NDRange so that all image data is covered by the parallelisation scheme. */ MitkOcl_EXPORT unsigned int iDivUp(unsigned int dividend, unsigned int divisor); -/*! \brief Loads a program source code from a text-file to a string, - -@param srcFilename: the file location -@param srcPreamble: will be added on the top (e.g. #define USE_UCHAR / USE_FLOAT) - -@return a char* -*/ -MitkOcl_EXPORT char* oclLoadProgramSource( const char* ,const char* ,size_t* ); - /** @brief Returns the name of an OpenCL Error as a string Most of the OpenCL Methods ( cl ) return an integer error code. This method translates the error value given as parameter to the corresponding error name. For example the value -30 will be translated to CL_INVALID_VALUE */ MitkOcl_EXPORT std::string GetOclErrorAsString( int _clErr ); /** @brief Checks whether the given value corresponds to an OpenCL Error value and prints this message out as MITK_ERROR if yes */ MitkOcl_EXPORT void GetOclError(int _clErr); /** @brief Returns a platform ID of an OpenCL-capable GPU, or throws an exception */ MitkOcl_EXPORT cl_int oclGetPlatformID(cl_platform_id* selectedPlatform); /*! \brief Prints out the essential support information about current device */ MitkOcl_EXPORT void oclPrintDeviceInfo(cl_device_id); /*! @brief Prints the available memory info about the given object to std::cout */ MitkOcl_EXPORT void oclPrintMemObjectInfo( cl_mem memobj); /*! \brief Checks the given code for errors and produces a std::cout output if the _err does not equal CL_SUCCESS. The output includes also the filename and the line number of the method call. */ MitkOcl_EXPORT bool oclCheckError(int _err, const char*, int); /*! \brief Logs the GPU Program binary code @param clProg: the OpenCL Program to log @param clDev: the OpenCL-capable device the program was tried to be compiled for */ MitkOcl_EXPORT void oclLogBinary(cl_program clProg, cl_device_id clDev); /*! \brief Shows the OpenCL-Program build info, called if clBuildProgram != CL_SUCCES @param clProg: the OpenCL Program to log @param clDev: the OpenCL-capable device the program was tried to be compiled for */ MitkOcl_EXPORT void oclLogBuildInfo(cl_program clProg, cl_device_id clDev); /** \brief Print out all supported image formats for given image type @param _type the image type ( CL_MEM_OBJECT_2D or CL_MEM_OBJECT_3D ) @param _context the OpenCL context to be examined */ MitkOcl_EXPORT void GetSupportedImageFormats(cl_context _context, cl_mem_object_type _type); /** @brief Translates the internal image type identifier to a human readable description string */ MitkOcl_EXPORT std::string GetImageTypeAsString( const unsigned int _in); #endif //mitkOclUtils_h