diff --git a/Modules/OpenCL/mitkOclFilter.cpp b/Modules/OpenCL/mitkOclFilter.cpp index 5312e59117..cdf84d053f 100644 --- a/Modules/OpenCL/mitkOclFilter.cpp +++ b/Modules/OpenCL/mitkOclFilter.cpp @@ -1,248 +1,248 @@ /*=================================================================== 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 //usService #include "usServiceReference.h" #include #include #include #include #include #include mitk::OclFilter::OclFilter() : m_ClCompilerFlags(""), m_ClProgram(nullptr), m_CommandQue(nullptr), m_FilterID("mitkOclFilter"), m_Preambel(" "), m_Initialized(false) { } mitk::OclFilter::OclFilter(const char* filename) : m_ClCompilerFlags(""), m_ClProgram(nullptr), m_CommandQue(nullptr), m_FilterID(filename), m_Preambel(" "), m_Initialized(false) { 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, nullptr, this->m_GlobalWorkSize, m_LocalWorkSize, 0, nullptr, nullptr); 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_ClFiles.empty()) { MITK_ERROR<<"No OpenCL Source FILE specified"; return false; } if (m_ClProgram == nullptr) { try { this->m_ClProgram = resources->GetProgram( this->m_FilterID ); } catch(const mitk::Exception& e) { - MITK_INFO << "Program not stored in resource manager, compiling."; + MITK_INFO << "Program not stored in resource manager, compiling. " << e; this->CompileSource(); } } return m_Initialized; } void mitk::OclFilter::LoadSourceFiles(CStringList &sourceCode, ClSizeList &sourceCodeSize) { for( CStringList::iterator it = m_ClFiles.begin(); it != m_ClFiles.end(); ++it ) { MITK_DEBUG << "Load 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); // read resource file to a string std::istreambuf_iterator eos; std::string source(std::istreambuf_iterator(rss), eos); // add preambel and build up string to compile std::string src(m_Preambel); src.append("\n"); src.append(source); // allocate new char buffer char* tmp = new char[src.size() + 1]; strcpy(tmp,src.c_str()); // add source to list sourceCode.push_back((const char*)tmp); sourceCodeSize.push_back(src.size()); } } void mitk::OclFilter::CompileSource() { // helper variable int clErr = 0; CStringList sourceCode; ClSizeList sourceCodeSize; if (m_ClFiles.empty()) { MITK_ERROR("ocl.filter") << "No shader source file was set"; return; } //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 LoadSourceFiles(sourceCode, sourceCodeSize); if ( !sourceCode.empty() ) { // create program from all files in the file list 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); MITK_DEBUG("ocl.filter") << "cl compiler flags: " << compilerOptions.c_str(); clErr = clBuildProgram(m_ClProgram, 0, nullptr, compilerOptions.c_str(), nullptr, nullptr); 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); // free the char buffers with the source code for( CStringList::iterator it = sourceCode.begin(); it != sourceCode.end(); ++it ) { delete[] *it; } } 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::AddSourceFile(const char* filename) { m_ClFiles.push_back(filename); } void mitk::OclFilter::SetCompilerFlags(const char* flags) { m_ClCompilerFlags = flags; } bool mitk::OclFilter::IsInitialized() { return m_Initialized; } diff --git a/Modules/OpenCL/mitkOclImageToImageFilter.cpp b/Modules/OpenCL/mitkOclImageToImageFilter.cpp index f78d6f9def..10f88b06a4 100644 --- a/Modules/OpenCL/mitkOclImageToImageFilter.cpp +++ b/Modules/OpenCL/mitkOclImageToImageFilter.cpp @@ -1,193 +1,193 @@ /*=================================================================== 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 "mitkOclImageToImageFilter.h" #include "mitkOclImage.h" #include "mitkException.h" mitk::OclImageToImageFilter::OclImageToImageFilter() { m_Output = mitk::OclImage::New(); } mitk::OclImageToImageFilter::~OclImageToImageFilter() { } mitk::OclImage::Pointer mitk::OclImageToImageFilter::GetGPUOutput() { // initialize some variables m_Output->SetPixelType(m_Input->GetPixelType()); // create new image, for passing the essential information to the output m_Output->InitializeMITKImage(); const unsigned int dimension = m_Input->GetDimension(); unsigned int* dimensions = m_Input->GetDimensions(); m_Output->SetDimensions( dimensions ); m_Output->SetDimension( (unsigned short)dimension ); m_Output->GetMITKImage()->Initialize( this->GetOutputType(), dimension, dimensions); const mitk::SlicedGeometry3D::Pointer p_slg = m_Input->GetMITKImage()->GetSlicedGeometry(0); m_Output->GetMITKImage()->SetSpacing( p_slg->GetSpacing() ); m_Output->GetMITKImage()->SetGeometry( m_Input->GetMITKImage()->GetGeometry() ); return this->m_Output; } mitk::Image::Pointer mitk::OclImageToImageFilter::GetOutput() { if (m_Output->IsModified(GPU_DATA)) { void* pData = m_Output->TransferDataToCPU(m_CommandQue); const unsigned int dimension = m_Input->GetDimension(); unsigned int* dimensions = m_Input->GetDimensions(); const mitk::SlicedGeometry3D::Pointer p_slg = m_Input->GetMITKImage()->GetSlicedGeometry(); MITK_DEBUG << "Creating new MITK Image."; m_Output->GetMITKImage()->Initialize( this->GetOutputType(), dimension, dimensions); m_Output->GetMITKImage()->SetSpacing( p_slg->GetSpacing()); m_Output->GetMITKImage()->SetGeometry( m_Input->GetMITKImage()->GetGeometry() ); m_Output->GetMITKImage()->SetImportVolume( pData, 0, 0, mitk::Image::ReferenceMemory); } MITK_DEBUG << "Image Initialized."; return m_Output->GetMITKImage(); } mitk::PixelType mitk::OclImageToImageFilter::GetOutputType() { // get the current image format from the input image const cl_image_format* currentImFormat = this->m_Input->GetPixelType(); // return the value according to the current channel type switch( currentImFormat->image_channel_data_type ) { case CL_UNORM_INT8: return mitk::MakeScalarPixelType(); case CL_UNSIGNED_INT8: return mitk::MakeScalarPixelType(); case CL_UNORM_INT16: return mitk::MakeScalarPixelType(); default: return mitk::MakeScalarPixelType(); } } int mitk::OclImageToImageFilter::GetBytesPerElem() { return (this->m_CurrentType + 1); } bool mitk::OclImageToImageFilter::InitExec(cl_kernel ckKernel) { cl_int clErr = 0; if( m_Input.IsNull() ) mitkThrow() << "Input image is null."; // get image size once const unsigned int uiImageWidth = m_Input->GetDimension(0); const unsigned int uiImageHeight = m_Input->GetDimension(1); const unsigned int uiImageDepth = m_Input->GetDimension(2); // compute work sizes this->SetWorkingSize( 8, uiImageWidth, 8, uiImageHeight , 8, uiImageDepth ); cl_mem clBuffIn = m_Input->GetGPUImage(this->m_CommandQue); cl_mem clBuffOut = m_Output->GetGPUBuffer(); if (!clBuffIn) { if ( m_Input->TransferDataToGPU(m_CommandQue) != CL_SUCCESS ) { mitkThrow()<< "Could not create / initialize gpu image."; } clBuffIn = m_Input->GetGPUImage(m_CommandQue); } // output image not initialized if (!clBuffOut) { //TODO bpp, or SetImageWidth/Height/... MITK_DEBUG << "Create GPU Image call " << uiImageWidth<< "x"<CreateGPUImage(uiImageWidth, uiImageHeight, uiImageDepth, this->m_CurrentType + 1); } clErr = 0; clErr = clSetKernelArg(ckKernel, 0, sizeof(cl_mem), &clBuffIn); clErr |= clSetKernelArg(ckKernel, 1, sizeof(cl_mem), &clBuffOut); CHECK_OCL_ERR( clErr ); if( clErr != CL_SUCCESS ) mitkThrow() << "OpenCL Part initialization failed with " << GetOclErrorAsString(clErr); return( clErr == CL_SUCCESS ); } bool mitk::OclImageToImageFilter::InitExec(cl_kernel ckKernel, unsigned int* dimensions) { cl_int clErr = 0; if( m_Input.IsNull() ) mitkThrow() << "Input image is null."; // get image size once const unsigned int uiImageWidth = dimensions[0]; const unsigned int uiImageHeight = dimensions[1]; - const unsigned int uiImageDepth = dimensions[2]; + const unsigned int uiImageDepth = dimensions[2]+1; // compute work sizes this->SetWorkingSize( 8, uiImageWidth, 8, uiImageHeight , 8, uiImageDepth ); cl_mem clBuffIn = m_Input->GetGPUImage(this->m_CommandQue); cl_mem clBuffOut = m_Output->GetGPUBuffer(); if (!clBuffIn) { if ( m_Input->TransferDataToGPU(m_CommandQue) != CL_SUCCESS ) { mitkThrow()<< "Could not create / initialize gpu image."; } clBuffIn = m_Input->GetGPUImage(m_CommandQue); } // output image not initialized //TODO bpp, or SetImageWidth/Height/... MITK_INFO << "Create GPU Image call " << uiImageWidth<< "x"<CreateGPUImage(uiImageWidth, uiImageHeight, uiImageDepth, this->m_CurrentType + 1); clErr = 0; clErr = clSetKernelArg(ckKernel, 0, sizeof(cl_mem), &clBuffIn); clErr |= clSetKernelArg(ckKernel, 1, sizeof(cl_mem), &clBuffOut); CHECK_OCL_ERR( clErr ); if( clErr != CL_SUCCESS ) mitkThrow() << "OpenCL Part initialization failed with " << GetOclErrorAsString(clErr); return( clErr == CL_SUCCESS ); } \ No newline at end of file diff --git a/Modules/PhotoacousticsAlgorithms/Algorithms/mitkPhotoacousticBeamformingFilter.cpp b/Modules/PhotoacousticsAlgorithms/Algorithms/mitkPhotoacousticBeamformingFilter.cpp index 47587fb719..ca03f169a8 100644 --- a/Modules/PhotoacousticsAlgorithms/Algorithms/mitkPhotoacousticBeamformingFilter.cpp +++ b/Modules/PhotoacousticsAlgorithms/Algorithms/mitkPhotoacousticBeamformingFilter.cpp @@ -1,607 +1,613 @@ /*=================================================================== mitkPhotoacousticBeamformingFilter 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 "mitkPhotoacousticBeamformingFilter.h" #include "mitkProperties.h" #include "mitkImageReadAccessor.h" #include #include #include #include #include #include #include "mitkImageCast.h" #include mitk::BeamformingFilter::BeamformingFilter() : m_OutputData(nullptr), m_InputData(nullptr) { this->SetNumberOfIndexedInputs(1); this->SetNumberOfRequiredInputs(1); m_ProgressHandle = [](int, std::string) {}; } void mitk::BeamformingFilter::SetProgressHandle(std::function progressHandle) { m_ProgressHandle = progressHandle; } mitk::BeamformingFilter::~BeamformingFilter() { } void mitk::BeamformingFilter::GenerateInputRequestedRegion() { Superclass::GenerateInputRequestedRegion(); mitk::Image* output = this->GetOutput(); mitk::Image* input = const_cast (this->GetInput()); if (!output->IsInitialized()) { return; } input->SetRequestedRegionToLargestPossibleRegion(); //GenerateTimeInInputRegion(output, input); } void mitk::BeamformingFilter::GenerateOutputInformation() { mitk::Image::ConstPointer input = this->GetInput(); mitk::Image::Pointer output = this->GetOutput(); if ((output->IsInitialized()) && (this->GetMTime() <= m_TimeOfHeaderInitialization.GetMTime())) return; itkDebugMacro(<< "GenerateOutputInformation()"); unsigned int dim[] = { m_Conf.ReconstructionLines, m_Conf.SamplesPerLine, input->GetDimension(2) }; output->Initialize(mitk::MakeScalarPixelType(), 3, dim); mitk::Vector3D spacing; spacing[0] = m_Conf.Pitch * m_Conf.TransducerElements * 1000 / m_Conf.ReconstructionLines; spacing[1] = m_Conf.RecordTime / 2 * m_Conf.SpeedOfSound * 1000 / m_Conf.SamplesPerLine; spacing[2] = 1; output->GetGeometry()->SetSpacing(spacing); output->GetGeometry()->Modified(); output->SetPropertyList(input->GetPropertyList()->Clone()); m_TimeOfHeaderInitialization.Modified(); } void mitk::BeamformingFilter::GenerateData() { GenerateOutputInformation(); mitk::Image::Pointer input = this->GetInput(); mitk::Image::Pointer output = this->GetOutput(); if (!output->IsInitialized()) return; float inputDim[2] = { (float)input->GetDimension(0), (float)input->GetDimension(1) }; float outputDim[2] = { (float)output->GetDimension(0), (float)output->GetDimension(1) }; unsigned short chunkSize = 2; // TODO: make this slightly less arbitrary unsigned int oclOutputDim[3] = { output->GetDimension(0), output->GetDimension(1), output->GetDimension(2) }; unsigned int oclOutputDimChunk[3] = { output->GetDimension(0), output->GetDimension(1), chunkSize}; unsigned int oclInputDimChunk[3] = { input->GetDimension(0), input->GetDimension(1), chunkSize}; unsigned int oclOutputDimLastChunk[3] = { output->GetDimension(0), output->GetDimension(1), input->GetDimension(2) % chunkSize }; unsigned int oclInputDimLastChunk[3] = { input->GetDimension(0), input->GetDimension(1), input->GetDimension(2) % chunkSize }; const int apodArraySize = m_Conf.TransducerElements * 4; // set the resolution of the apodization array float* ApodWindow; // calculate the appropiate apodization window switch (m_Conf.Apod) { case beamformingSettings::Apodization::Hann: ApodWindow = VonHannFunction(apodArraySize); break; case beamformingSettings::Apodization::Hamm: ApodWindow = HammFunction(apodArraySize); break; case beamformingSettings::Apodization::Box: ApodWindow = BoxFunction(apodArraySize); break; default: ApodWindow = BoxFunction(apodArraySize); break; } int progInterval = output->GetDimension(2) / 20 > 1 ? output->GetDimension(2) / 20 : 1; // the interval at which we update the gui progress bar auto begin = std::chrono::high_resolution_clock::now(); // debbuging the performance... if (!m_Conf.UseGPU) { for (unsigned int i = 0; i < output->GetDimension(2); ++i) // seperate Slices should get Beamforming seperately applied { mitk::ImageReadAccessor inputReadAccessor(input, input->GetSliceData(i)); m_OutputData = new float[m_Conf.ReconstructionLines*m_Conf.SamplesPerLine]; m_InputDataPuffer = new float[input->GetDimension(0)*input->GetDimension(1)]; // first, we convert any data to float, which we use by default if (input->GetPixelType().GetTypeAsString() == "scalar (float)" || input->GetPixelType().GetTypeAsString() == " (float)") { m_InputData = (float*)inputReadAccessor.GetData(); } else { MITK_INFO << "Pixel type is not float, abort"; return; } // fill the image with zeros for (int l = 0; l < outputDim[0]; ++l) { for (int s = 0; s < outputDim[1]; ++s) { m_OutputData[l*(short)outputDim[1] + s] = 0; } } std::thread *threads = new std::thread[(short)outputDim[0]]; // every line will be beamformed in a seperate thread if (m_Conf.Algorithm == beamformingSettings::BeamformingAlgorithm::DAS) { if (m_Conf.DelayCalculationMethod == beamformingSettings::DelayCalc::QuadApprox) { for (short line = 0; line < outputDim[0]; ++line) { threads[line] = std::thread(&BeamformingFilter::DASQuadraticLine, this, m_InputData, m_OutputData, inputDim, outputDim, line, ApodWindow, apodArraySize); } } else if (m_Conf.DelayCalculationMethod == beamformingSettings::DelayCalc::Spherical) { for (short line = 0; line < outputDim[0]; ++line) { threads[line] = std::thread(&BeamformingFilter::DASSphericalLine, this, m_InputData, m_OutputData, inputDim, outputDim, line, ApodWindow, apodArraySize); } } } else if (m_Conf.Algorithm == beamformingSettings::BeamformingAlgorithm::DMAS) { if (m_Conf.DelayCalculationMethod == beamformingSettings::DelayCalc::QuadApprox) { for (short line = 0; line < outputDim[0]; ++line) { threads[line] = std::thread(&BeamformingFilter::DMASQuadraticLine, this, m_InputData, m_OutputData, inputDim, outputDim, line, ApodWindow, apodArraySize); } } else if (m_Conf.DelayCalculationMethod == beamformingSettings::DelayCalc::Spherical) { for (short line = 0; line < outputDim[0]; ++line) { threads[line] = std::thread(&BeamformingFilter::DMASSphericalLine, this, m_InputData, m_OutputData, inputDim, outputDim, line, ApodWindow, apodArraySize); } } } // wait for all lines to finish for (short line = 0; line < outputDim[0]; ++line) { threads[line].join(); } output->SetSlice(m_OutputData, i); if (i % progInterval == 0) m_ProgressHandle((int)((i + 1) / (float)output->GetDimension(2) * 100), "performing reconstruction"); delete[] m_OutputData; delete[] m_InputDataPuffer; m_OutputData = nullptr; m_InputData = nullptr; } } else { mitk::PhotoacousticOCLBeamformer::Pointer m_oclFilter = mitk::PhotoacousticOCLBeamformer::New(); try { if (m_Conf.Algorithm == beamformingSettings::BeamformingAlgorithm::DAS) { if (m_Conf.DelayCalculationMethod == beamformingSettings::DelayCalc::QuadApprox) m_oclFilter->SetAlgorithm(PhotoacousticOCLBeamformer::BeamformingAlgorithm::DASQuad, true); else if (m_Conf.DelayCalculationMethod == beamformingSettings::DelayCalc::Spherical) m_oclFilter->SetAlgorithm(PhotoacousticOCLBeamformer::BeamformingAlgorithm::DASSphe, true); } else if (m_Conf.Algorithm == beamformingSettings::BeamformingAlgorithm::DMAS) { if (m_Conf.DelayCalculationMethod == beamformingSettings::DelayCalc::QuadApprox) m_oclFilter->SetAlgorithm(PhotoacousticOCLBeamformer::BeamformingAlgorithm::DMASQuad, true); else if (m_Conf.DelayCalculationMethod == beamformingSettings::DelayCalc::Spherical) m_oclFilter->SetAlgorithm(PhotoacousticOCLBeamformer::BeamformingAlgorithm::DMASSphe, true); } m_oclFilter->SetApodisation(ApodWindow, apodArraySize); m_oclFilter->SetOutputDim(oclOutputDimChunk); m_oclFilter->SetBeamformingParameters(m_Conf.SpeedOfSound, m_Conf.TimeSpacing, m_Conf.Pitch, m_Conf.Angle, m_Conf.Photoacoustic, m_Conf.TransducerElements); - + if (chunkSize < oclOutputDim[2]) { bool skip = false; for (unsigned int i = 0; !skip && i < ceil((float)oclOutputDim[2] / (float)chunkSize); ++i) { m_ProgressHandle(100 * ((float)(i * chunkSize) / (float)oclOutputDim[2]), "performing reconstruction"); mitk::Image::Pointer chunk = mitk::Image::New(); if ((int)((oclOutputDim[2]) - (i * chunkSize)) == (int)(1 + chunkSize)) { // A 3d image of 3rd dimension == 1 can not be processed by openCL, make sure that this case never arises oclInputDimLastChunk[2] = input->GetDimension(2) % chunkSize + chunkSize; oclOutputDimLastChunk[2] = input->GetDimension(2) % chunkSize + chunkSize; chunk->Initialize(input->GetPixelType(), 3, oclInputDimLastChunk); m_oclFilter->SetOutputDim(oclOutputDimLastChunk); skip = true; //skip the last chunk } else if ((oclOutputDim[2]) - (i * chunkSize) >= chunkSize) chunk->Initialize(input->GetPixelType(), 3, oclInputDimChunk); else { chunk->Initialize(input->GetPixelType(), 3, oclInputDimLastChunk); m_oclFilter->SetOutputDim(oclOutputDimLastChunk); } chunk->SetSpacing(input->GetGeometry()->GetSpacing()); mitk::ImageReadAccessor ChunkMove(input); chunk->SetImportVolume((void*)&(((float*)const_cast(ChunkMove.GetData()))[i * chunkSize * input->GetDimension(0) * input->GetDimension(1)]), 0, 0, mitk::Image::ReferenceMemory); m_oclFilter->SetInput(chunk); m_oclFilter->Update(); auto out = m_oclFilter->GetOutput(); + + unsigned int lastCopiedSlice = 0; + if(skip) + { + lastCopiedSlice = i * chunkSize + 1; + } for (unsigned int s = i * chunkSize; s < oclOutputDim[2]; ++s) // TODO: make the bounds here smaller... { mitk::ImageReadAccessor copy(out, out->GetSliceData(s - i * chunkSize)); output->SetImportSlice(const_cast(copy.GetData()), s, 0, 0, mitk::Image::ReferenceMemory); } } } else { m_ProgressHandle(50, "performing reconstruction"); m_oclFilter->SetOutputDim(oclOutputDim); m_oclFilter->SetInput(input); m_oclFilter->Update(); auto out = m_oclFilter->GetOutput(); mitk::ImageReadAccessor copy(out); output->SetImportVolume(const_cast(copy.GetData()), 0, 0, mitk::Image::ReferenceMemory); } } catch (mitk::Exception &e) { std::string errorMessage = "Caught unexpected exception "; errorMessage.append(e.what()); MITK_ERROR << errorMessage; } } m_TimeOfHeaderInitialization.Modified(); auto end = std::chrono::high_resolution_clock::now(); MITK_INFO << "Beamforming of " << output->GetDimension(2) << " Images completed in " << ((float)std::chrono::duration_cast(end - begin).count()) / 1000000 << "ms" << std::endl; } void mitk::BeamformingFilter::Configure(beamformingSettings settings) { m_Conf = settings; } float* mitk::BeamformingFilter::VonHannFunction(int samples) { float* ApodWindow = new float[samples]; for (int n = 0; n < samples; ++n) { ApodWindow[n] = (1 - cos(2 * M_PI * n / (samples - 1))) / 2; } return ApodWindow; } float* mitk::BeamformingFilter::HammFunction(int samples) { float* ApodWindow = new float[samples]; for (int n = 0; n < samples; ++n) { ApodWindow[n] = 0.54 - 0.46*cos(2 * M_PI*n / (samples - 1)); } return ApodWindow; } float* mitk::BeamformingFilter::BoxFunction(int samples) { float* ApodWindow = new float[samples]; for (int n = 0; n < samples; ++n) { ApodWindow[n] = 1; } return ApodWindow; } void mitk::BeamformingFilter::DASQuadraticLine(float* input, float* output, float inputDim[2], float outputDim[2], const short& line, float* apodisation, const short& apodArraySize) { float& inputS = inputDim[1]; float& inputL = inputDim[0]; float& outputS = outputDim[1]; float& outputL = outputDim[0]; short AddSample = 0; short maxLine = 0; short minLine = 0; float delayMultiplicator = 0; float l_i = 0; float s_i = 0; float part = 0.07 * inputL; float tan_phi = std::tan(m_Conf.Angle / 360 * 2 * M_PI); float part_multiplicator = tan_phi * m_Conf.TimeSpacing * m_Conf.SpeedOfSound / m_Conf.Pitch * m_Conf.ReconstructionLines / m_Conf.TransducerElements; float apod_mult = 1; short usedLines = (maxLine - minLine); //quadratic delay l_i = line / outputL * inputL; for (short sample = 0; sample < outputS; ++sample) { s_i = (float)sample / outputS * inputS / 2; part = part_multiplicator*s_i; if (part < 1) part = 1; maxLine = (short)std::min((l_i + part) + 1, inputL); minLine = (short)std::max((l_i - part), 0.0f); usedLines = (maxLine - minLine); apod_mult = apodArraySize / (maxLine - minLine); delayMultiplicator = pow((1 / (m_Conf.TimeSpacing*m_Conf.SpeedOfSound) * (m_Conf.Pitch*m_Conf.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 - m_Conf.Photoacoustic)*s_i; if (AddSample < inputS && AddSample >= 0) output[sample*(short)outputL + line] += input[l_s + AddSample*(short)inputL] * apodisation[(short)((l_s - minLine)*apod_mult)]; else --usedLines; } output[sample*(short)outputL + line] = output[sample*(short)outputL + line] / usedLines; } } void mitk::BeamformingFilter::DASSphericalLine(float* input, float* output, float inputDim[2], float outputDim[2], const short& line, float* apodisation, const short& apodArraySize) { float& inputS = inputDim[1]; float& inputL = inputDim[0]; float& outputS = outputDim[1]; float& outputL = outputDim[0]; short AddSample = 0; short maxLine = 0; short minLine = 0; float l_i = 0; float s_i = 0; float part = 0.07 * inputL; float tan_phi = std::tan(m_Conf.Angle / 360 * 2 * M_PI); float part_multiplicator = tan_phi * m_Conf.TimeSpacing * m_Conf.SpeedOfSound / m_Conf.Pitch * m_Conf.ReconstructionLines / m_Conf.TransducerElements; float apod_mult = 1; short usedLines = (maxLine - minLine); //exact delay l_i = (float)line / outputL * inputL; for (short sample = 0; sample < outputS; ++sample) { s_i = (float)sample / outputS * inputS / 2; part = part_multiplicator*s_i; if (part < 1) part = 1; maxLine = (short)std::min((l_i + part) + 1, inputL); minLine = (short)std::max((l_i - part), 0.0f); usedLines = (maxLine - minLine); apod_mult = apodArraySize / (maxLine - minLine); for (short l_s = minLine; l_s < maxLine; ++l_s) { AddSample = (int)sqrt( pow(s_i, 2) + pow((1 / (m_Conf.TimeSpacing*m_Conf.SpeedOfSound) * ((l_s - l_i)*m_Conf.Pitch*m_Conf.TransducerElements) / inputL), 2) ) + (1 - m_Conf.Photoacoustic)*s_i; if (AddSample < inputS && AddSample >= 0) output[sample*(short)outputL + line] += input[l_s + AddSample*(short)inputL] * apodisation[(short)((l_s - minLine)*apod_mult)]; else --usedLines; } output[sample*(short)outputL + line] = output[sample*(short)outputL + line] / usedLines; } } void mitk::BeamformingFilter::DMASQuadraticLine(float* input, float* output, float inputDim[2], float outputDim[2], const short& line, float* apodisation, const short& apodArraySize) { float& inputS = inputDim[1]; float& inputL = inputDim[0]; float& outputS = outputDim[1]; float& outputL = outputDim[0]; short maxLine = 0; short minLine = 0; float delayMultiplicator = 0; float l_i = 0; float s_i = 0; float part = 0.07 * inputL; float tan_phi = std::tan(m_Conf.Angle / 360 * 2 * M_PI); float part_multiplicator = tan_phi * m_Conf.TimeSpacing * m_Conf.SpeedOfSound / m_Conf.Pitch * m_Conf.ReconstructionLines / m_Conf.TransducerElements; float apod_mult = 1; float mult = 0; short usedLines = (maxLine - minLine); //quadratic delay l_i = line / outputL * inputL; for (short sample = 0; sample < outputS; ++sample) { s_i = sample / outputS * inputS / 2; part = part_multiplicator*s_i; if (part < 1) part = 1; maxLine = (short)std::min((l_i + part) + 1, inputL); minLine = (short)std::max((l_i - part), 0.0f); usedLines = (maxLine - minLine); apod_mult = apodArraySize / (maxLine - minLine); delayMultiplicator = pow((1 / (m_Conf.TimeSpacing*m_Conf.SpeedOfSound) * (m_Conf.Pitch*m_Conf.TransducerElements) / inputL), 2) / s_i / 2; //calculate the AddSamples beforehand to save some time short* AddSample = new short[maxLine - minLine]; for (short l_s = 0; l_s < maxLine - minLine; ++l_s) { AddSample[l_s] = (short)(delayMultiplicator * pow((minLine + l_s - l_i), 2) + s_i) + (1 - m_Conf.Photoacoustic)*s_i; } for (short l_s1 = minLine; l_s1 < maxLine - 1; ++l_s1) { if (AddSample[l_s1 - minLine] < (short)inputS && AddSample[l_s1 - minLine] >= 0) { for (short l_s2 = l_s1 + 1; l_s2 < maxLine; ++l_s2) { if (AddSample[l_s2 - minLine] < inputS && AddSample[l_s2 - minLine] >= 0) { mult = input[l_s2 + AddSample[l_s2 - minLine] * (short)inputL] * apodisation[(short)((l_s2 - minLine)*apod_mult)] * input[l_s1 + AddSample[l_s1 - minLine] * (short)inputL] * apodisation[(short)((l_s1 - minLine)*apod_mult)]; output[sample*(short)outputL + line] += sqrt(abs(mult)) * ((mult > 0) - (mult < 0)); } } } else --usedLines; } output[sample*(short)outputL + line] = 10 * output[sample*(short)outputL + line] / (pow(usedLines, 2) - (usedLines - 1)); delete[] AddSample; } } void mitk::BeamformingFilter::DMASSphericalLine(float* input, float* output, float inputDim[2], float outputDim[2], const short& line, float* apodisation, const short& apodArraySize) { float& inputS = inputDim[1]; float& inputL = inputDim[0]; float& outputS = outputDim[1]; float& outputL = outputDim[0]; short maxLine = 0; short minLine = 0; float l_i = 0; float s_i = 0; float part = 0.07 * inputL; float tan_phi = std::tan(m_Conf.Angle / 360 * 2 * M_PI); float part_multiplicator = tan_phi * m_Conf.TimeSpacing * m_Conf.SpeedOfSound / m_Conf.Pitch * m_Conf.ReconstructionLines / m_Conf.TransducerElements; float apod_mult = 1; float mult = 0; short usedLines = (maxLine - minLine); //exact delay l_i = line / outputL * inputL; for (short sample = 0; sample < outputS; ++sample) { s_i = sample / outputS * inputS / 2; part = part_multiplicator*s_i; if (part < 1) part = 1; maxLine = (short)std::min((l_i + part) + 1, inputL); minLine = (short)std::max((l_i - part), 0.0f); usedLines = (maxLine - minLine); apod_mult = apodArraySize / (maxLine - minLine); //calculate the AddSamples beforehand to save some time short* AddSample = new short[maxLine - minLine]; for (short l_s = 0; l_s < maxLine - minLine; ++l_s) { AddSample[l_s] = (short)sqrt( pow(s_i, 2) + pow((1 / (m_Conf.TimeSpacing*m_Conf.SpeedOfSound) * ((minLine + l_s - l_i)*m_Conf.Pitch*m_Conf.TransducerElements) / inputL), 2) ) + (1 - m_Conf.Photoacoustic)*s_i; } for (short l_s1 = minLine; l_s1 < maxLine - 1; ++l_s1) { if (AddSample[l_s1 - minLine] < inputS && AddSample[l_s1 - minLine] >= 0) { for (short l_s2 = l_s1 + 1; l_s2 < maxLine; ++l_s2) { if (AddSample[l_s2 - minLine] < inputS && AddSample[l_s2 - minLine] >= 0) { mult = input[l_s2 + AddSample[l_s2 - minLine] * (short)inputL] * apodisation[(int)((l_s2 - minLine)*apod_mult)] * input[l_s1 + AddSample[l_s1 - minLine] * (short)inputL] * apodisation[(int)((l_s1 - minLine)*apod_mult)]; output[sample*(short)outputL + line] += sqrt(abs(mult)) * ((mult > 0) - (mult < 0)); } } } else --usedLines; } output[sample*(short)outputL + line] = 10 * output[sample*(short)outputL + line] / (pow(usedLines, 2) - (usedLines - 1)); delete[] AddSample; } }