diff --git a/Modules/PhotoacousticsAlgorithms/Resources/DASQuadratic.cl b/Modules/PhotoacousticsAlgorithms/Resources/DASQuadratic.cl deleted file mode 100644 index 12a1456807..0000000000 --- a/Modules/PhotoacousticsAlgorithms/Resources/DASQuadratic.cl +++ /dev/null @@ -1,72 +0,0 @@ -/*=================================================================== - -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 ckDASQuad( - __global float* dSource, // input image - __global float* dDest, // output buffer - __global float* apodArray, - unsigned short apodArraySize, - float SpeedOfSound, - float TimeSpacing, - float Pitch, - float Angle, - unsigned short PAImage, - unsigned short TransducerElements, - unsigned int inputL, - unsigned int inputS, - unsigned int Slices, - unsigned int outputL, - unsigned int outputS // 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); - - // terminate non-valid threads - if ( globalPosX < outputL && globalPosY < outputS && globalPosZ < Slices ) - { - float l_i = (float)globalPosX / outputL * inputL; - float s_i = (float)globalPosY / outputS * inputS / 2; - - float part = (tan(Angle / 360 * 2 * M_PI) * ((SpeedOfSound * TimeSpacing) * s_i) / (Pitch * TransducerElements)) * inputL; - if (part < 1) - part = 1; - - short maxLine = min((l_i + part) + 1, (float)inputL); - short minLine = max((l_i - part), 0.0f); - short usedLines = (maxLine - minLine); - float apod_mult = (float)apodArraySize / (float)usedLines; - - short AddSample = 0; - float output = 0; - - float delayMultiplicator = pow(1 / (TimeSpacing*SpeedOfSound) * Pitch * TransducerElements / inputL, 2) / s_i / 2; - - for (short l_s = minLine; l_s < maxLine; ++l_s) - { - AddSample = delayMultiplicator * pow((l_s - l_i), 2) + s_i + (1-PAImage)*s_i; - if (AddSample < inputS && AddSample >= 0) { - output += apodArray[(int)((l_s - minLine)*apod_mult)] * dSource[(int)(globalPosZ * inputL * inputS + AddSample * inputL + l_s)]; - } - else - --usedLines; - } - - dDest[ globalPosZ * outputL * outputS + globalPosY * outputL + globalPosX ] = output / usedLines; - } -} \ No newline at end of file diff --git a/Modules/PhotoacousticsAlgorithms/Resources/DASSpherical.cl b/Modules/PhotoacousticsAlgorithms/Resources/DASSpherical.cl deleted file mode 100644 index 7cedb732a0..0000000000 --- a/Modules/PhotoacousticsAlgorithms/Resources/DASSpherical.cl +++ /dev/null @@ -1,73 +0,0 @@ -/*=================================================================== - -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 ckDASSphe( - __global float* dSource, // input image - __global float* dDest, // output buffer - __global float* apodArray, - unsigned short apodArraySize, - float SpeedOfSound, - float TimeSpacing, - float Pitch, - float Angle, - unsigned short PAImage, - unsigned short TransducerElements, - unsigned int inputL, - unsigned int inputS, - unsigned int Slices, - unsigned int outputL, - unsigned int outputS // 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); - - // terminate non-valid threads - if ( globalPosX < outputL && globalPosY < outputS && globalPosZ < Slices ) - { - float l_i = (float)globalPosX / outputL * inputL; - float s_i = (float)globalPosY / outputS * inputS / 2; - - float part = (tan(Angle / 360 * 2 * M_PI) * ((SpeedOfSound * TimeSpacing) * s_i) / (Pitch * TransducerElements)) * inputL; - if (part < 1) - part = 1; - - short maxLine = min((l_i + part) + 1, (float)inputL); - short minLine = max((l_i - part), 0.0f); - short usedLines = (maxLine - minLine); - float apod_mult = (float)apodArraySize / (float)usedLines; - - short AddSample = 0; - float output = 0; - - for (short l_s = minLine; l_s < maxLine; ++l_s) - { - AddSample = sqrt( - pow(s_i, 2) - + - pow((1 / (TimeSpacing*SpeedOfSound) * ((l_s - l_i)*Pitch*TransducerElements) / inputL), 2) - ) + (1-PAImage)*s_i; - if (AddSample < inputS && AddSample >= 0) - output += apodArray[(short)((l_s - minLine)*apod_mult)] * dSource[(unsigned int)(globalPosZ * inputL * inputS + AddSample * inputL + l_s)]; - else - --usedLines; - } - - dDest[ globalPosZ * outputL * outputS + globalPosY * outputL + globalPosX ] = output / usedLines; - } -} \ No newline at end of file diff --git a/Modules/PhotoacousticsAlgorithms/Resources/DMASQuadratic.cl b/Modules/PhotoacousticsAlgorithms/Resources/DMASQuadratic.cl deleted file mode 100644 index ec7ec8a868..0000000000 --- a/Modules/PhotoacousticsAlgorithms/Resources/DMASQuadratic.cl +++ /dev/null @@ -1,84 +0,0 @@ -/*=================================================================== - -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 ckDMASQuad( - __global float* dSource, // input image - __global float* dDest, // output buffer - __global float* apodArray, - unsigned short apodArraySize, - float SpeedOfSound, - float TimeSpacing, - float Pitch, - float Angle, - unsigned short PAImage, - unsigned short TransducerElements, - unsigned int inputL, - unsigned int inputS, - unsigned int Slices, - unsigned int outputL, - unsigned int outputS // 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); - - // terminate non-valid threads - if ( globalPosX < outputL && globalPosY < outputS && globalPosZ < Slices ) - { - float l_i = (float)globalPosX / outputL * inputL; - float s_i = (float)globalPosY / outputS * inputS / 2; - - float part = (tan(Angle / 360 * 2 * M_PI) * ((SpeedOfSound * TimeSpacing) * s_i) / (Pitch * TransducerElements)) * inputL; - if (part < 1) - part = 1; - - short maxLine = min((l_i + part) + 1, (float)inputL); - short minLine = max((l_i - part), 0.0f); - short usedLines = (maxLine - minLine); - float apod_mult = apodArraySize / (maxLine - minLine); - - short AddSample1 = 0; - short AddSample2 = 0; - - float output = 0; - float mult = 0; - - float delayMultiplicator = pow(1 / (TimeSpacing*SpeedOfSound) * Pitch * TransducerElements / inputL, 2) / s_i / 2; - - for (short l_s1 = minLine; l_s1 < maxLine; ++l_s1) - { - AddSample1 = delayMultiplicator * pow((l_s1 - l_i), 2) + s_i + (1-PAImage)*s_i; - if (AddSample1 < inputS && AddSample1 >= 0) { - for (short l_s2 = l_s1 + 1; l_s2 < maxLine; ++l_s2) - { - AddSample2 = delayMultiplicator * pow((l_s2 - l_i), 2) + s_i + (1-PAImage)*s_i; - if (AddSample1 < inputS && AddSample1 >= 0) { - mult = apodArray[(short)((l_s2 - minLine)*apod_mult)] * dSource[(unsigned int)(globalPosZ * inputL * inputS + AddSample2 * inputL + l_s2)] - * apodArray[(short)((l_s1 - minLine)*apod_mult)] * dSource[(unsigned int)(globalPosZ * inputL * inputS + AddSample1 * inputL + l_s1)]; - output += sqrt(mult * ((float)(mult>0)-(float)(mult<0))) * ((mult > 0) - (mult < 0)); - - } - } - } - else - --usedLines; - } - - dDest[ globalPosZ * outputL * outputS + globalPosY * outputL + globalPosX ] = output / (pow((float)usedLines, 2.0f) - (usedLines - 1)); - } -} \ No newline at end of file diff --git a/Modules/PhotoacousticsAlgorithms/Resources/DMASSpherical.cl b/Modules/PhotoacousticsAlgorithms/Resources/DMASSpherical.cl deleted file mode 100644 index a883160ead..0000000000 --- a/Modules/PhotoacousticsAlgorithms/Resources/DMASSpherical.cl +++ /dev/null @@ -1,94 +0,0 @@ -/*=================================================================== - -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 ckDMASSphe( - __global float* dSource, // input image - __global float* dDest, // output buffer - __global float* apodArray, - unsigned short apodArraySize, - float SpeedOfSound, - float TimeSpacing, - float Pitch, - float Angle, - unsigned short PAImage, - unsigned short TransducerElements, - unsigned int inputL, - unsigned int inputS, - unsigned int Slices, - unsigned int outputL, - unsigned int outputS // 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); - - // terminate non-valid threads - if ( globalPosX < outputL && globalPosY < outputS && globalPosZ < Slices ) - { - float l_i = (float)globalPosX / outputL * inputL; - float s_i = (float)globalPosY / outputS * inputS / 2; - - float part = (tan(Angle / 360 * 2 * M_PI) * ((SpeedOfSound * TimeSpacing) * s_i) / (Pitch * TransducerElements)) * inputL; - if (part < 1) - part = 1; - - short maxLine = min((l_i + part) + 1, (float)inputL); - short minLine = max((l_i - part), 0.0f); - short usedLines = (maxLine - minLine); - float apod_mult = apodArraySize / (maxLine - minLine); - - short AddSample1 = 0; - short AddSample2 = 0; - - float output = 0; - float mult = 0; - - float delayMultiplicator = pow(1 / (TimeSpacing*SpeedOfSound) * Pitch * TransducerElements / inputL, 2) / s_i / 2; - - for (short l_s1 = minLine; l_s1 < maxLine; ++l_s1) - { - AddSample1 = sqrt( - pow(s_i, 2) - + - pow((1 / (TimeSpacing*SpeedOfSound) * ((l_s1 - l_i)*Pitch*TransducerElements)/inputL), 2) - ) + (1-PAImage)*s_i; - if (AddSample1 < inputS && AddSample1 >= 0) { - for (short l_s2 = l_s1 + 1; l_s2 < maxLine; ++l_s2) - { - AddSample2 = sqrt( - pow(s_i, 2) - + - pow((1 / (TimeSpacing*SpeedOfSound) * ((l_s2 - l_i)*Pitch*TransducerElements)/inputL), 2) - ) + (1-PAImage)*s_i; - if (AddSample1 < inputS && AddSample1 >= 0) { - mult = apodArray[(short)((l_s2 - minLine)*apod_mult)] * dSource[(unsigned int)(globalPosZ * inputL * inputS + AddSample2 * inputL + l_s2)] - * apodArray[(short)((l_s1 - minLine)*apod_mult)] * dSource[(unsigned int)(globalPosZ * inputL * inputS + AddSample1 * inputL + l_s1)]; - output += sqrt(mult * ((float)(mult>0)-(float)(mult<0))) * ((mult > 0) - (mult < 0)); - - } - } - } - else - --usedLines; - } - - dDest[ globalPosZ * outputL * outputS + globalPosY * outputL + globalPosX ] = output / (pow((float)usedLines, 2.0f) - (usedLines - 1)); - } -} - - \ No newline at end of file diff --git a/Modules/PhotoacousticsAlgorithms/Resources/MemoryLocSum.cl b/Modules/PhotoacousticsAlgorithms/Resources/MemoryLocSum.cl deleted file mode 100644 index 4bc9b3bfa3..0000000000 --- a/Modules/PhotoacousticsAlgorithms/Resources/MemoryLocSum.cl +++ /dev/null @@ -1,35 +0,0 @@ -/*=================================================================== - -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 ckMemoryLocSum( __global unsigned short *input, - __global unsigned int *sums, - __global unsigned int *finalSum - ) - { - uint id = get_global_id(0) + get_global_id(1) * get_global_size(0); - - uint sum = 0; - - for (uint pos = 0; pos < id; ++pos) - { - sum += input[3 * pos]; - } - - sums[id] = sum; - - if (id == (get_global_size(0) - 1) + (get_global_size(1) - 1) * get_global_size(0)) - finalSum[0] = sum + input[3 * id]; - } \ No newline at end of file diff --git a/Modules/PhotoacousticsAlgorithms/files.cmake b/Modules/PhotoacousticsAlgorithms/files.cmake index fc025f27fb..88f4bd6e46 100644 --- a/Modules/PhotoacousticsAlgorithms/files.cmake +++ b/Modules/PhotoacousticsAlgorithms/files.cmake @@ -1,23 +1,19 @@ file(GLOB_RECURSE H_FILES RELATIVE "${CMAKE_CURRENT_SOURCE_DIR}" "${CMAKE_CURRENT_SOURCE_DIR}/include/*") set(CPP_FILES source/mitkPhotoacousticImage.cpp source/mitkPhotoacousticBeamformingFilter.cpp source/OpenCLFilter/mitkPhotoacousticOCLBeamformingFilter.cpp source/OpenCLFilter/mitkPhotoacousticBModeFilter.cpp source/OpenCLFilter/mitkPhotoacousticOCLUsedLinesCalculation.cpp source/OpenCLFilter/mitkPhotoacousticOCLDelayCalculation.cpp - source/OpenCLFilter/mitkPhotoacousticOCLMemoryLocSum.cpp ) set(RESOURCE_FILES - DASQuadratic.cl - DASSpherical.cl BModeAbs.cl BModeAbsLog.cl UsedLinesCalculation.cl - MemoryLocSum.cl DelayCalculation.cl DMAS.cl DAS.cl ) diff --git a/Modules/PhotoacousticsAlgorithms/include/OpenCLFilter/mitkPhotoacousticOCLBeamformingFilter.h b/Modules/PhotoacousticsAlgorithms/include/OpenCLFilter/mitkPhotoacousticOCLBeamformingFilter.h index 060ee5c0e3..2f7c33d971 100644 --- a/Modules/PhotoacousticsAlgorithms/include/OpenCLFilter/mitkPhotoacousticOCLBeamformingFilter.h +++ b/Modules/PhotoacousticsAlgorithms/include/OpenCLFilter/mitkPhotoacousticOCLBeamformingFilter.h @@ -1,140 +1,138 @@ /*=================================================================== 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 _MITKPHOTOACOUSTICSOCLBEAMFORMER_H_ #define _MITKPHOTOACOUSTICSOCLBEAMFORMER_H_ #ifdef PHOTOACOUSTICS_USE_GPU #include "mitkOclDataSetToDataSetFilter.h" #include #include "mitkPhotoacousticOCLDelayCalculation.h" -#include "mitkPhotoacousticOCLMemoryLocSum.h" #include "mitkPhotoacousticOCLUsedLinesCalculation.h" #include "mitkPhotoacousticBeamformingSettings.h" #include namespace mitk { /** 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 PhotoacousticOCLBeamformingFilter : public OclDataSetToDataSetFilter, public itk::Object { public: mitkClassMacroItkParent(PhotoacousticOCLBeamformingFilter, itk::Object); itkNewMacro(Self); /** * @brief SetInput Set the input data through an image. Arbitrary images are supported */ void SetInput(Image::Pointer image); /** * @brief SetInput Manually set the input data while providing dimensions and memory size of the input data. */ void SetInput(void* data, unsigned int* dimensions, unsigned int BpE); void* GetOutput(); /** * @brief GetOutputAsImage Returns an mitk::Image constructed from the processed data */ mitk::Image::Pointer GetOutputAsImage(); /** Update the filter */ void Update(); /** Set the Apodisation function to apply when beamforming */ void SetApodisation(float* apodisation, unsigned short apodArraySize) { m_ApodArraySize = apodArraySize; m_Apodisation = apodisation; } void SetConfig(BeamformingSettings settings) { m_ConfOld = m_Conf; m_Conf = settings; } protected: /** Constructor */ PhotoacousticOCLBeamformingFilter(); /** Destructor */ virtual ~PhotoacousticOCLBeamformingFilter(); /** Initialize the filter */ bool Initialize(); void UpdateDataBuffers(); void Execute(); mitk::PixelType GetOutputType() { return mitk::MakeScalarPixelType(); } int GetBytesPerElem() { return sizeof(float); } virtual us::Module* GetModule(); private: /** The OpenCL kernel for the filter */ cl_kernel m_PixelCalculation; unsigned int m_OutputDim[3]; float* m_Apodisation; unsigned short m_ApodArraySize; unsigned short m_PAImage; BeamformingSettings m_Conf; BeamformingSettings m_ConfOld; mitk::Image::Pointer m_InputImage; size_t m_ChunkSize[3]; - mitk::OCLMemoryLocSum::Pointer m_SumFilter; mitk::OCLUsedLinesCalculation::Pointer m_UsedLinesCalculation; mitk::OCLDelayCalculation::Pointer m_DelayCalculation; cl_mem m_ApodizationBuffer; cl_mem m_MemoryLocationsBuffer; cl_mem m_DelaysBuffer; cl_mem m_UsedLinesBuffer; std::chrono::steady_clock::time_point m_Begin; }; } #endif #endif diff --git a/Modules/PhotoacousticsAlgorithms/include/OpenCLFilter/mitkPhotoacousticOCLMemoryLocSum.h b/Modules/PhotoacousticsAlgorithms/include/OpenCLFilter/mitkPhotoacousticOCLMemoryLocSum.h deleted file mode 100644 index bce56edbc2..0000000000 --- a/Modules/PhotoacousticsAlgorithms/include/OpenCLFilter/mitkPhotoacousticOCLMemoryLocSum.h +++ /dev/null @@ -1,99 +0,0 @@ -/*=================================================================== - -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 _MITKPHOTOACOUSTICSPARALLELSUM_H_ -#define _MITKPHOTOACOUSTICSPARALLELSUM_H_ - -#ifdef PHOTOACOUSTICS_USE_GPU - -#include "mitkOclDataSetToDataSetFilter.h" -#include - -namespace mitk -{ - /** 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 OCLMemoryLocSum : public OclDataSetToDataSetFilter, public itk::Object - { - - public: - mitkClassMacroItkParent(OCLMemoryLocSum, itk::Object); - itkNewMacro(Self); - - /** - * @brief SetInput Set the input image. Only 3D images are supported for now. - * @param image a 3D image. - * @throw mitk::Exception if the dimesion is not 3. - */ - - /** Update the filter */ - void Update(); - - void SetInputDimensions(unsigned int* dimensions) - { - m_Dim[0] = dimensions[0]; - m_Dim[1] = dimensions[1]; - } - - unsigned int GetSum() - { - return m_Sum; - } - - protected: - - /** Constructor */ - OCLMemoryLocSum(); - - /** Destructor */ - virtual ~OCLMemoryLocSum(); - - /** 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; - - unsigned int m_Sum; - unsigned int m_Dim[3]; - }; -} -#endif -#endif \ No newline at end of file diff --git a/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLBeamformingFilter.cpp b/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLBeamformingFilter.cpp index b529bea54a..72187362aa 100644 --- a/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLBeamformingFilter.cpp +++ b/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLBeamformingFilter.cpp @@ -1,250 +1,249 @@ /*=================================================================== 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. ===================================================================*/ #ifdef PHOTOACOUSTICS_USE_GPU #include "./OpenCLFilter/mitkPhotoacousticOCLBeamformingFilter.h" #include "usServiceReference.h" mitk::PhotoacousticOCLBeamformingFilter::PhotoacousticOCLBeamformingFilter() : m_PixelCalculation( NULL ), m_InputImage(mitk::Image::New()), m_ApodizationBuffer(nullptr), m_MemoryLocationsBuffer(nullptr), m_DelaysBuffer(nullptr), m_UsedLinesBuffer(nullptr) { - this->AddSourceFile("DASQuadratic.cl"); this->AddSourceFile("DAS.cl"); this->AddSourceFile("DMAS.cl"); this->m_FilterID = "OpenCLBeamformingFilter"; this->Initialize(); unsigned int dim[] = { 128, 2048, 2 }; mitk::Vector3D spacing; spacing[0] = 1; spacing[1] = 1; spacing[2] = 1; m_InputImage->Initialize(mitk::MakeScalarPixelType(), 3, dim); m_InputImage->SetSpacing(spacing); m_ChunkSize[0] = 128; m_ChunkSize[1] = 128; m_ChunkSize[2] = 8; m_UsedLinesCalculation = mitk::OCLUsedLinesCalculation::New(); m_DelayCalculation = mitk::OCLDelayCalculation::New(); } mitk::PhotoacousticOCLBeamformingFilter::~PhotoacousticOCLBeamformingFilter() { if ( this->m_PixelCalculation ) { clReleaseKernel( m_PixelCalculation ); } if (m_ApodizationBuffer) clReleaseMemObject(m_ApodizationBuffer); } void mitk::PhotoacousticOCLBeamformingFilter::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::PhotoacousticOCLBeamformingFilter::UpdateDataBuffers() { //Initialize the Output try { MITK_INFO << "Updating Workgroup size for new dimensions"; size_t outputSize = m_Conf.ReconstructionLines * m_Conf.SamplesPerLine * m_Conf.inputDim[2]; m_OutputDim[0] = m_Conf.ReconstructionLines; m_OutputDim[1] = m_Conf.SamplesPerLine; m_OutputDim[2] = m_Conf.inputDim[2]; this->InitExec(this->m_PixelCalculation, m_OutputDim, outputSize, sizeof(float)); } catch (const mitk::Exception& e) { MITK_ERROR << "Caught exception while initializing filter: " << e.what(); return; } if (BeamformingSettings::SettingsChangedOpenCL(m_Conf, m_ConfOld)) { cl_int clErr = 0; MITK_INFO << "Updating Buffers for new configuration"; // create the apodisation buffer if (m_Apodisation == nullptr) { MITK_INFO << "No apodisation function set; Beamforming will be done without any apodisation."; m_Apodisation = new float[1]; m_Apodisation[0] = 1; m_ApodArraySize = 1; } us::ServiceReference ref = GetModuleContext()->GetServiceReference(); OclResourceService* resources = GetModuleContext()->GetService(ref); cl_context gpuContext = resources->GetContext(); if (m_ApodizationBuffer) clReleaseMemObject(m_ApodizationBuffer); this->m_ApodizationBuffer = clCreateBuffer(gpuContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(float) * m_ApodArraySize, m_Apodisation, &clErr); CHECK_OCL_ERR(clErr); // calculate used lines m_UsedLinesCalculation->SetConfig(m_Conf); m_UsedLinesCalculation->Update(); m_UsedLinesBuffer = m_UsedLinesCalculation->GetGPUOutput()->GetGPUBuffer(); // calculate the Delays m_DelayCalculation->SetConfig(m_Conf); m_DelayCalculation->SetInputs(m_UsedLinesBuffer); m_DelayCalculation->Update(); m_DelaysBuffer = m_DelayCalculation->GetGPUOutput()->GetGPUBuffer(); m_ConfOld = m_Conf; } } void mitk::PhotoacousticOCLBeamformingFilter::Execute() { cl_int clErr = 0; UpdateDataBuffers(); clErr = clSetKernelArg(this->m_PixelCalculation, 2, sizeof(cl_mem), &(this->m_UsedLinesBuffer)); clErr |= clSetKernelArg(this->m_PixelCalculation, 3, sizeof(cl_mem), &(this->m_DelaysBuffer)); clErr |= clSetKernelArg(this->m_PixelCalculation, 4, sizeof(cl_mem), &(this->m_ApodizationBuffer)); clErr |= clSetKernelArg(this->m_PixelCalculation, 5, sizeof(cl_ushort), &(this->m_ApodArraySize)); clErr |= clSetKernelArg(this->m_PixelCalculation, 6, sizeof(cl_uint), &(this->m_Conf.inputDim[0])); clErr |= clSetKernelArg(this->m_PixelCalculation, 7, sizeof(cl_uint), &(this->m_Conf.inputDim[1])); clErr |= clSetKernelArg(this->m_PixelCalculation, 8, sizeof(cl_uint), &(this->m_Conf.inputDim[2])); clErr |= clSetKernelArg(this->m_PixelCalculation, 9, sizeof(cl_uint), &(this->m_Conf.ReconstructionLines)); clErr |= clSetKernelArg(this->m_PixelCalculation, 10, sizeof(cl_uint), &(this->m_Conf.SamplesPerLine)); // execute the filter on a 3D NDRange if (m_OutputDim[2] == 1 || m_ChunkSize[2] == 1) this->ExecuteKernelChunks(m_PixelCalculation, 2, m_ChunkSize); else this->ExecuteKernelChunks(m_PixelCalculation, 3, m_ChunkSize); // signalize the GPU-side data changed m_Output->Modified( GPU_DATA ); } us::Module *mitk::PhotoacousticOCLBeamformingFilter::GetModule() { return us::GetModuleContext()->GetModule(); } bool mitk::PhotoacousticOCLBeamformingFilter::Initialize() { bool buildErr = true; cl_int clErr = 0; if ( OclFilter::Initialize() ) { switch (m_Conf.Algorithm) { case BeamformingSettings::BeamformingAlgorithm::DAS: { this->m_PixelCalculation = clCreateKernel(this->m_ClProgram, "ckDAS", &clErr); break; } case BeamformingSettings::BeamformingAlgorithm::DMAS: { this->m_PixelCalculation = clCreateKernel(this->m_ClProgram, "ckDMAS", &clErr); break; } default: { MITK_INFO << "No beamforming algorithm specified, setting to DAS"; this->m_PixelCalculation = clCreateKernel(this->m_ClProgram, "ckDAS", &clErr); break; } } buildErr |= CHECK_OCL_ERR( clErr ); } CHECK_OCL_ERR(clErr); return (OclFilter::IsInitialized() && buildErr ); } void mitk::PhotoacousticOCLBeamformingFilter::SetInput(mitk::Image::Pointer image) { OclDataSetToDataSetFilter::SetInput(image); m_InputImage = image; m_Conf.inputDim[0] = m_InputImage->GetDimension(0); m_Conf.inputDim[1] = m_InputImage->GetDimension(1); m_Conf.inputDim[2] = m_InputImage->GetDimension(2); } void mitk::PhotoacousticOCLBeamformingFilter::SetInput(void* data, unsigned int* dimensions, unsigned int BpE) { OclDataSetToDataSetFilter::SetInput(data, dimensions[0] * dimensions[1] * dimensions[2], BpE); m_Conf.inputDim[0] = dimensions[0]; m_Conf.inputDim[1] = dimensions[1]; m_Conf.inputDim[2] = dimensions[2]; } mitk::Image::Pointer mitk::PhotoacousticOCLBeamformingFilter::GetOutputAsImage() { mitk::Image::Pointer outputImage = mitk::Image::New(); if (m_Output->IsModified(GPU_DATA)) { void* pData = m_Output->TransferDataToCPU(m_CommandQue); const unsigned int dimension = 3; unsigned int dimensions[3] = { (unsigned int)m_OutputDim[0], (unsigned int)m_OutputDim[1], (unsigned int)m_OutputDim[2] }; const mitk::SlicedGeometry3D::Pointer p_slg = m_InputImage->GetSlicedGeometry(); MITK_DEBUG << "Creating new MITK Image."; outputImage->Initialize(this->GetOutputType(), dimension, dimensions); outputImage->SetSpacing(p_slg->GetSpacing()); outputImage->SetImportVolume(pData, 0, 0, mitk::Image::ReferenceMemory); } MITK_DEBUG << "Image Initialized."; return outputImage; } void* mitk::PhotoacousticOCLBeamformingFilter::GetOutput() { return OclDataSetToDataSetFilter::GetOutput(); } #endif diff --git a/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLMemoryLocSum.cpp b/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLMemoryLocSum.cpp deleted file mode 100644 index ebfac4ec3f..0000000000 --- a/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLMemoryLocSum.cpp +++ /dev/null @@ -1,119 +0,0 @@ -/*=================================================================== - -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. - -===================================================================*/ - -#define _USE_MATH_DEFINES - -#include -#include "./OpenCLFilter/mitkPhotoacousticOCLMemoryLocSum.h" -#include "usServiceReference.h" -#include "mitkImageReadAccessor.h" - -mitk::OCLMemoryLocSum::OCLMemoryLocSum() - : m_PixelCalculation(NULL), m_Sum(0) -{ - this->AddSourceFile("MemoryLocSum.cl"); - this->m_FilterID = "MemoryLocSum"; - - this->Initialize(); -} - -mitk::OCLMemoryLocSum::~OCLMemoryLocSum() -{ - if (this->m_PixelCalculation) - { - clReleaseKernel(m_PixelCalculation); - } -} - -void mitk::OCLMemoryLocSum::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::OCLMemoryLocSum::Execute() -{ - cl_int clErr = 0; - - unsigned int gridDim[3] = { m_Dim[0], m_Dim[1], 1 }; - size_t outputSize = gridDim[0] * gridDim[1] * 1; - - try - { - this->InitExec(this->m_PixelCalculation, gridDim, outputSize, sizeof(unsigned int)); - } - catch (const mitk::Exception& e) - { - MITK_ERROR << "Caught exception while initializing filter: " << e.what(); - return; - } - - us::ServiceReference ref = GetModuleContext()->GetServiceReference(); - OclResourceService* resources = GetModuleContext()->GetService(ref); - cl_context gpuContext = resources->GetContext(); - - cl_mem cl_Sum = clCreateBuffer(gpuContext, CL_MEM_WRITE_ONLY, sizeof(unsigned int), NULL, &clErr); - CHECK_OCL_ERR(clErr); - - clErr = clSetKernelArg(this->m_PixelCalculation, 2, sizeof(cl_mem), &cl_Sum); - - CHECK_OCL_ERR(clErr); - - // execute the filter on a 3D NDRange - this->ExecuteKernel(m_PixelCalculation, 2); - - // signalize the GPU-side data changed - m_Output->Modified(GPU_DATA); - - char* data = new char[1 * sizeof(unsigned int)]; - - clErr = clEnqueueReadBuffer(m_CommandQue, cl_Sum, CL_FALSE, 0, 1 * sizeof(unsigned int), data, 0, nullptr, nullptr); - CHECK_OCL_ERR(clErr); - - clFlush(m_CommandQue); - - m_Sum = ((unsigned int*)data)[0]; -} - -us::Module *mitk::OCLMemoryLocSum::GetModule() -{ - return us::GetModuleContext()->GetModule(); -} - -bool mitk::OCLMemoryLocSum::Initialize() -{ - bool buildErr = true; - cl_int clErr = 0; - - if (OclFilter::Initialize()) - { - this->m_PixelCalculation = clCreateKernel(this->m_ClProgram, "ckMemoryLocSum", &clErr); - buildErr |= CHECK_OCL_ERR(clErr); - } - return (OclFilter::IsInitialized() && buildErr); -} \ No newline at end of file