diff --git a/Modules/PhotoacousticsAlgorithms/Resources/UsedLinesCalculation.cl b/Modules/PhotoacousticsAlgorithms/Resources/UsedLinesCalculation.cl index a0d4af6fb7..638de2b279 100644 --- a/Modules/PhotoacousticsAlgorithms/Resources/UsedLinesCalculation.cl +++ b/Modules/PhotoacousticsAlgorithms/Resources/UsedLinesCalculation.cl @@ -1,48 +1,50 @@ /*=================================================================== 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 ckUsedLines( __global unsigned short* dDest, // output buffer float partMult, unsigned int inputL, - unsigned int inputS + unsigned int inputS, + unsigned int outputL, + unsigned int outputS ) { // get thread identifier unsigned int globalPosX = get_global_id(0); unsigned int globalPosY = get_global_id(1); - unsigned short outputS = get_global_size(1); - unsigned short outputL = get_global_size(0); + //unsigned short outputS = get_global_size(1); + //unsigned short outputL = get_global_size(0); // terminate non-valid threads if ( globalPosX < outputL && globalPosY < outputS) { float l_i = (float)globalPosX / outputL * inputL; float s_i = (float)globalPosY / outputS * inputS / 2; float part = partMult * s_i; if (part < 1) part = 1; unsigned short maxLine = min((l_i + part) + 1, (float)inputL); unsigned short minLine = max((l_i - part), 0.0f); dDest[globalPosY * 3 * outputL + 3 * globalPosX] = (maxLine - minLine); //usedLines dDest[globalPosY * 3 * outputL + 3 * globalPosX + 1] = minLine; //minLine dDest[globalPosY * 3 * outputL + 3 * globalPosX + 2] = maxLine; //maxLine } } \ No newline at end of file diff --git a/Modules/PhotoacousticsAlgorithms/include/OpenCLFilter/mitkPhotoacousticOCLUsedLinesCalculation.h b/Modules/PhotoacousticsAlgorithms/include/OpenCLFilter/mitkPhotoacousticOCLUsedLinesCalculation.h index 6ee3cf0ccf..e787d2009a 100644 --- a/Modules/PhotoacousticsAlgorithms/include/OpenCLFilter/mitkPhotoacousticOCLUsedLinesCalculation.h +++ b/Modules/PhotoacousticsAlgorithms/include/OpenCLFilter/mitkPhotoacousticOCLUsedLinesCalculation.h @@ -1,89 +1,90 @@ /*=================================================================== 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 _MITKPHOTOACOUSTICSOCLUSEDLINESCALCULATION_H_ #define _MITKPHOTOACOUSTICSOCLUSEDLINESCALCULATION_H_ #if defined(PHOTOACOUSTICS_USE_GPU) || DOXYGEN #include "mitkOclDataSetToDataSetFilter.h" #include #include "mitkPhotoacousticBeamformingSettings.h" namespace mitk { /*! * \brief Class implementing a mitk::OclDataSetToDataSetFilter to calculate which lines each sample should use when beamforming. * * The class must be given a configuration class instance of mitk::BeamformingSettings for beamforming parameters through mitk::OCLDelayCalculation::SetConfig(BeamformingSettings conf) */ class OCLUsedLinesCalculation : public OclDataSetToDataSetFilter, public itk::Object { public: mitkClassMacroItkParent(OCLUsedLinesCalculation, itk::Object); itkNewMacro(Self); void Update(); /** \brief Sets a new configuration to use. * * @param conf The configuration set to use for the calculation of the used lines. */ void SetConfig(BeamformingSettings conf) { m_Conf = conf; } protected: /** Constructor */ OCLUsedLinesCalculation(); /** Destructor */ virtual ~OCLUsedLinesCalculation(); /** Initialize the filter */ bool Initialize(); void Execute(); mitk::PixelType GetOutputType() { return mitk::MakeScalarPixelType(); } int GetBytesPerElem() { return sizeof(unsigned short); } virtual us::Module* GetModule(); int m_sizeThis; private: /** The OpenCL kernel for the filter */ cl_kernel m_PixelCalculation; BeamformingSettings m_Conf; float m_part; + size_t m_ChunkSize[3]; }; } #endif #endif diff --git a/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLUsedLinesCalculation.cpp b/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLUsedLinesCalculation.cpp index 349c4781c4..7a03d85a10 100644 --- a/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLUsedLinesCalculation.cpp +++ b/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLUsedLinesCalculation.cpp @@ -1,108 +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. ===================================================================*/ #if defined(PHOTOACOUSTICS_USE_GPU) || DOXYGEN #define _USE_MATH_DEFINES #include #include "./OpenCLFilter/mitkPhotoacousticOCLUsedLinesCalculation.h" #include "usServiceReference.h" #include "mitkImageReadAccessor.h" mitk::OCLUsedLinesCalculation::OCLUsedLinesCalculation() : m_PixelCalculation(NULL) { this->AddSourceFile("UsedLinesCalculation.cl"); this->m_FilterID = "UsedLinesCalculation"; + m_ChunkSize[0] = 128; + m_ChunkSize[1] = 128; + m_ChunkSize[2] = 8; + this->Initialize(); } mitk::OCLUsedLinesCalculation::~OCLUsedLinesCalculation() { if (this->m_PixelCalculation) { clReleaseKernel(m_PixelCalculation); } } void mitk::OCLUsedLinesCalculation::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::OCLUsedLinesCalculation::Execute() { cl_int clErr = 0; unsigned int gridDim[3] = { m_Conf.ReconstructionLines, m_Conf.SamplesPerLine, 1 }; size_t outputSize = gridDim[0] * gridDim[1] * 3; try { this->InitExecNoInput(this->m_PixelCalculation, gridDim, outputSize, sizeof(unsigned short)); } catch (const mitk::Exception& e) { MITK_ERROR << "Caught exception while initializing filter: " << e.what(); return; } m_part = (tan(m_Conf.Angle / 360 * 2 * M_PI) * ((m_Conf.SpeedOfSound * m_Conf.TimeSpacing)) / (m_Conf.Pitch * m_Conf.TransducerElements)) * m_Conf.inputDim[0]; clErr = clSetKernelArg(this->m_PixelCalculation, 1, sizeof(cl_float), &(this->m_part)); clErr |= clSetKernelArg(this->m_PixelCalculation, 2, sizeof(cl_uint), &(this->m_Conf.inputDim[0])); clErr |= clSetKernelArg(this->m_PixelCalculation, 3, sizeof(cl_uint), &(this->m_Conf.inputDim[1])); + clErr |= clSetKernelArg(this->m_PixelCalculation, 4, sizeof(cl_uint), &(this->m_Conf.ReconstructionLines)); + clErr |= clSetKernelArg(this->m_PixelCalculation, 5, sizeof(cl_uint), &(this->m_Conf.SamplesPerLine)); CHECK_OCL_ERR(clErr); - // execute the filter on a 3D NDRange - this->ExecuteKernel(m_PixelCalculation, 2); + // execute the filter on a 2D NDRange + if (!this->ExecuteKernelChunks(m_PixelCalculation, 2, m_ChunkSize)) + mitkThrow() << "openCL Error when executing Kernel"; // signalize the GPU-side data changed m_Output->Modified(GPU_DATA); } us::Module *mitk::OCLUsedLinesCalculation::GetModule() { return us::GetModuleContext()->GetModule(); } bool mitk::OCLUsedLinesCalculation::Initialize() { bool buildErr = true; cl_int clErr = 0; if (OclFilter::Initialize()) { this->m_PixelCalculation = clCreateKernel(this->m_ClProgram, "ckUsedLines", &clErr); buildErr |= CHECK_OCL_ERR(clErr); } return (OclFilter::IsInitialized() && buildErr); } #endif