diff --git a/Modules/OpenCL/mitkOclFilter.cpp b/Modules/OpenCL/mitkOclFilter.cpp index 7965a24484..98d8902987 100644 --- a/Modules/OpenCL/mitkOclFilter.cpp +++ b/Modules/OpenCL/mitkOclFilter.cpp @@ -1,347 +1,348 @@ /*=================================================================== 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 +//standard library +#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 ); } -#include - bool mitk::OclFilter::ExecuteKernelChunks( cl_kernel kernel, unsigned int workSizeDim, size_t* chunksDim ) { size_t offset[3] ={0, 0, 0}; cl_int clErr = 0; if(workSizeDim == 2) { for(offset[0] = 0; offset[0] < m_GlobalWorkSize[0]; offset[0] += chunksDim[0]) { for(offset[1] = 0; offset[1] < m_GlobalWorkSize[1]; offset[1] += chunksDim[1]) { clErr |= clEnqueueNDRangeKernel(this->m_CommandQue, kernel, workSizeDim, offset, chunksDim, m_LocalWorkSize, 0, nullptr, nullptr); } } } else if(workSizeDim == 3) { for(offset[0] = 0; offset[0] < m_GlobalWorkSize[0]; offset[0] += chunksDim[0]) { for(offset[1] = 0; offset[1] < m_GlobalWorkSize[1]; offset[1] += chunksDim[1]) { for(offset[2] = 0; offset[2] < m_GlobalWorkSize[2]; offset[2] += chunksDim[2]) { clErr |= clEnqueueNDRangeKernel( this->m_CommandQue, kernel, workSizeDim, offset, chunksDim, m_LocalWorkSize, 0, nullptr, nullptr); } } } } CHECK_OCL_ERR(clErr); return ( clErr == CL_SUCCESS ); } bool mitk::OclFilter::ExecuteKernelChunksInBatches(cl_kernel kernel, unsigned int workSizeDim, size_t* chunksDim, size_t batchSize, int waitTimems) { size_t offset[3] = { 0, 0, 0 }; cl_int clErr = 0; unsigned int currentChunk = 0; cl_event* waitFor = new cl_event[batchSize]; if (workSizeDim == 2) { for (offset[0] = 0; offset[0] < m_GlobalWorkSize[0]; offset[0] += chunksDim[0]) { for (offset[1] = 0; offset[1] < m_GlobalWorkSize[1]; offset[1] += chunksDim[1]) { if (currentChunk % batchSize == 0 && currentChunk != 0) { clWaitForEvents(batchSize, &waitFor[0]); std::this_thread::sleep_for(std::chrono::milliseconds(waitTimems)); clErr |= clEnqueueNDRangeKernel(this->m_CommandQue, kernel, workSizeDim, offset, chunksDim, m_LocalWorkSize, 0, nullptr, &waitFor[0]); } else { clErr |= clEnqueueNDRangeKernel(this->m_CommandQue, kernel, workSizeDim, offset, chunksDim, m_LocalWorkSize, 0, nullptr, &waitFor[currentChunk % batchSize]); } currentChunk++; } } } else if (workSizeDim == 3) { for (offset[0] = 0; offset[0] < m_GlobalWorkSize[0]; offset[0] += chunksDim[0]) { for (offset[1] = 0; offset[1] < m_GlobalWorkSize[1]; offset[1] += chunksDim[1]) { for (offset[2] = 0; offset[2] < m_GlobalWorkSize[2]; offset[2] += chunksDim[2]) { if (currentChunk % batchSize == 0 && currentChunk != 0) { clWaitForEvents(batchSize, &waitFor[0]); std::this_thread::sleep_for(std::chrono::milliseconds(waitTimems)); clErr |= clEnqueueNDRangeKernel(this->m_CommandQue, kernel, workSizeDim, offset, chunksDim, m_LocalWorkSize, 0, nullptr, &waitFor[0]); } else { clErr |= clEnqueueNDRangeKernel(this->m_CommandQue, kernel, workSizeDim, offset, chunksDim, m_LocalWorkSize, 0, nullptr, &waitFor[currentChunk % batchSize]); } currentChunk++; } } } } - 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. " << 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/PhotoacousticsAlgorithms/Resources/DAS.cl b/Modules/PhotoacousticsAlgorithms/Resources/DAS.cl index 8a4c95135b..2540ec58db 100644 --- a/Modules/PhotoacousticsAlgorithms/Resources/DAS.cl +++ b/Modules/PhotoacousticsAlgorithms/Resources/DAS.cl @@ -1,64 +1,64 @@ /*=================================================================== 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 ckDAS( __global float* dSource, // input image __global float* dDest, // output buffer __global unsigned short* usedLines, __global unsigned short* delays, __constant float* apodArray, unsigned short apodArraySize, 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 / (float)outputL * (float)inputL; unsigned short curUsedLines = usedLines[globalPosY * 3 * outputL + 3 * globalPosX]; unsigned short minLine = usedLines[globalPosY * 3 * outputL + 3 * globalPosX + 1]; unsigned short maxLine = usedLines[globalPosY * 3 *outputL + 3 * globalPosX + 2]; float apod_mult = (float)apodArraySize / (float)curUsedLines; unsigned short Delay = 0; float output = 0; float mult = 0; for (short l_s = minLine; l_s < maxLine; ++l_s) { - Delay = delays[globalPosY * outputL + (int)(fabs(l_s - l_i)/(float)inputL * (float)outputL)]; + Delay = delays[globalPosY * (outputL / 2) + (int)(fabs(l_s - l_i)/(float)inputL * (float)outputL)]; if (Delay < inputS && Delay >= 0) { output += apodArray[(int)((l_s - minLine)*apod_mult)] * dSource[(int)(globalPosZ * inputL * inputS + Delay * inputL + l_s)]; } else --curUsedLines; } dDest[ globalPosZ * outputL * outputS + globalPosY * outputL + globalPosX ] = output / (float)curUsedLines; } } \ No newline at end of file diff --git a/Modules/PhotoacousticsAlgorithms/Resources/DMAS.cl b/Modules/PhotoacousticsAlgorithms/Resources/DMAS.cl index 7ec678b6ac..89cc58bf12 100644 --- a/Modules/PhotoacousticsAlgorithms/Resources/DMAS.cl +++ b/Modules/PhotoacousticsAlgorithms/Resources/DMAS.cl @@ -1,85 +1,85 @@ /*=================================================================== 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 ckDMAS( __global float* dSource, // input image __global float* dDest, // output buffer __global unsigned short* usedLines, __global unsigned short* AddSamples, __constant float* apodArray, unsigned short apodArraySize, 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 / (float)outputL * (float)inputL; unsigned short curUsedLines = usedLines[globalPosY * 3 * outputL + 3 * globalPosX]; unsigned short minLine = usedLines[globalPosY * 3 * outputL + 3 * globalPosX + 1]; unsigned short maxLine = usedLines[globalPosY * 3 *outputL + 3 * globalPosX + 2]; float apod_mult = (float)apodArraySize / (float)curUsedLines; unsigned short Delay1 = 0; unsigned short Delay2 = 0; float output = 0; float mult = 0; float s_1 = 0; float s_2 = 0; float apod_1 = 0; for (short l_s1 = minLine; l_s1 < maxLine; ++l_s1) { - Delay1 = AddSamples[globalPosY * outputL + (int)(fabs(l_s1 - l_i)/(float)inputL * (float)outputL)]; + Delay1 = AddSamples[globalPosY * (outputL / 2) + (int)(fabs(l_s1 - l_i)/(float)inputL * (float)outputL)]; if (Delay1 < inputS && Delay1 >= 0) { s_1 = dSource[(int)(globalPosZ * inputL * inputS + Delay1 * inputL + l_s1)]; apod_1 = apodArray[(int)((l_s1 - minLine)*apod_mult)]; for (short l_s2 = l_s1 + 1; l_s2 < maxLine; ++l_s2) { - Delay2 = AddSamples[globalPosY * outputL + (int)(fabs(l_s2 - l_i)/(float)inputL * (float)outputL)]; + Delay2 = AddSamples[globalPosY * (outputL / 2) + (int)(fabs(l_s2 - l_i)/(float)inputL * (float)outputL)]; if (Delay2 < inputS && Delay2 >= 0) { s_2 = dSource[(int)(globalPosZ * inputL * inputS + Delay2 * inputL + l_s2)]; mult = apodArray[(int)((l_s2 - minLine)*apod_mult)] * s_2 * apod_1 * s_1; output += sqrt(fabs(mult)) * ((mult > 0) - (mult < 0)); } } } else --curUsedLines; } dDest[ globalPosZ * outputL * outputS + globalPosY * outputL + globalPosX ] = output / (float)(curUsedLines * curUsedLines - (curUsedLines - 1)); } } diff --git a/Modules/PhotoacousticsAlgorithms/Resources/DelayCalculation.cl b/Modules/PhotoacousticsAlgorithms/Resources/DelayCalculation.cl index a388e23dc2..0be471b2b4 100644 --- a/Modules/PhotoacousticsAlgorithms/Resources/DelayCalculation.cl +++ b/Modules/PhotoacousticsAlgorithms/Resources/DelayCalculation.cl @@ -1,69 +1,69 @@ /*=================================================================== 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 ckDelayCalculationQuad( __global unsigned short *gDest, __global unsigned short *usedLines, unsigned int inputL, unsigned int inputS, unsigned int outputL, unsigned int outputS, char isPAImage, float delayMultiplicatorRaw // parameters ) { uint globalPosX = get_global_id(0); uint globalPosY = get_global_id(1); if (globalPosX * 2 < outputL && globalPosY < outputS) { float l_i = 0; // we calculate the delays relative to line zero float s_i = (float)globalPosY / (float)outputS * (float)inputS / 2; float l_s = (float)globalPosX / (float)outputL * (float)inputL; // the currently calculated line float delayMultiplicator = delayMultiplicatorRaw / s_i; - gDest[globalPosY * outputL + globalPosX] = delayMultiplicator * pow((l_s - l_i), 2) + s_i + (1-isPAImage)*s_i; + gDest[globalPosY * (outputL / 2) + globalPosX] = delayMultiplicator * pow((l_s - l_i), 2) + s_i + (1-isPAImage)*s_i; } } __kernel void ckDelayCalculationSphe( __global unsigned short *gDest, __global unsigned short *usedLines, unsigned int inputL, unsigned int inputS, unsigned int outputL, unsigned int outputS, char isPAImage, float delayMultiplicatorRaw // parameters ) { uint globalPosX = get_global_id(0); uint globalPosY = get_global_id(1); - if (globalPosX * 2< outputL && globalPosY < outputS) + if (globalPosX * 2 < outputL && globalPosY < outputS) { float l_i = 0; // we calculate the delays relative to line zero float s_i = (float)globalPosY / (float)outputS * (float)inputS / 2; float l_s = (float)globalPosX / (float)outputL * (float)inputL; // the currently calculated line - gDest[globalPosY * outputL + globalPosX] = + gDest[globalPosY * (outputL / 2) + globalPosX] = sqrt( pow(s_i, 2) + pow((delayMultiplicatorRaw * ((l_s - l_i)) / inputL), 2) ) + (1-isPAImage)*s_i; } } \ No newline at end of file diff --git a/Modules/PhotoacousticsAlgorithms/Resources/sDMAS.cl b/Modules/PhotoacousticsAlgorithms/Resources/sDMAS.cl index 1a8ff0269e..b362047846 100644 --- a/Modules/PhotoacousticsAlgorithms/Resources/sDMAS.cl +++ b/Modules/PhotoacousticsAlgorithms/Resources/sDMAS.cl @@ -1,87 +1,87 @@ /*=================================================================== 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 cksDMAS( __global float* dSource, // input image __global float* dDest, // output buffer __global unsigned short* usedLines, __global unsigned short* AddSamples, __constant float* apodArray, unsigned short apodArraySize, 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 / (float)outputL * (float)inputL; unsigned short curUsedLines = usedLines[globalPosY * 3 * outputL + 3 * globalPosX]; unsigned short minLine = usedLines[globalPosY * 3 * outputL + 3 * globalPosX + 1]; unsigned short maxLine = usedLines[globalPosY * 3 *outputL + 3 * globalPosX + 2]; float apod_mult = (float)apodArraySize / (float)curUsedLines; unsigned short Delay1 = 0; unsigned short Delay2 = 0; float output = 0; float mult = 0; float s_1 = 0; float s_2 = 0; float sign = 0; float apod_1 = 0; for (short l_s1 = minLine; l_s1 < maxLine; ++l_s1) { - Delay1 = AddSamples[globalPosY * outputL + (int)(fabs(l_s1 - l_i)/(float)inputL * (float)outputL)]; + Delay1 = AddSamples[globalPosY * (outputL / 2) + (int)(fabs(l_s1 - l_i)/(float)inputL * (float)outputL)]; if (Delay1 < inputS && Delay1 >= 0) { s_1 = dSource[(int)(globalPosZ * inputL * inputS + Delay1 * inputL + l_s1)]; apod_1 = apodArray[(int)((l_s1 - minLine)*apod_mult)]; sign += s_1; for (short l_s2 = l_s1 + 1; l_s2 < maxLine; ++l_s2) { - Delay2 = AddSamples[globalPosY * outputL + (int)(fabs(l_s2 - l_i)/(float)inputL * (float)outputL)]; + Delay2 = AddSamples[globalPosY * (outputL / 2) + (int)(fabs(l_s2 - l_i)/(float)inputL * (float)outputL)]; if (Delay2 < inputS && Delay2 >= 0) { s_2 = dSource[(int)(globalPosZ * inputL * inputS + Delay2 * inputL + l_s2)]; mult = apodArray[(int)((l_s2 - minLine)*apod_mult)] * s_2 * apod_1 * s_1; output += sqrt(fabs(mult)) * ((mult > 0) - (mult < 0)); } } } else --curUsedLines; } dDest[ globalPosZ * outputL * outputS + globalPosY * outputL + globalPosX ] = output / (float)(curUsedLines * curUsedLines - (curUsedLines - 1)) * ((sign > 0) - (sign < 0)); } } diff --git a/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLDelayCalculation.cpp b/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLDelayCalculation.cpp index 53019b8b57..a6730652d0 100644 --- a/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLDelayCalculation.cpp +++ b/Modules/PhotoacousticsAlgorithms/source/OpenCLFilter/mitkPhotoacousticOCLDelayCalculation.cpp @@ -1,124 +1,123 @@ /*=================================================================== 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/mitkPhotoacousticOCLDelayCalculation.h" #include "usServiceReference.h" #include "mitkImageReadAccessor.h" mitk::OCLDelayCalculation::OCLDelayCalculation() : m_PixelCalculation(NULL) { this->AddSourceFile("DelayCalculation.cl"); this->m_FilterID = "DelayCalculation"; m_ChunkSize[0] = 128; m_ChunkSize[1] = 128; m_ChunkSize[2] = 8; this->Initialize(); } mitk::OCLDelayCalculation::~OCLDelayCalculation() { if (this->m_PixelCalculation) { clReleaseKernel(m_PixelCalculation); } } void mitk::OCLDelayCalculation::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::OCLDelayCalculation::Execute() { cl_int clErr = 0; unsigned int gridDim[3] = { m_Conf.ReconstructionLines / 2, m_Conf.SamplesPerLine, 1 }; m_BufferSize = gridDim[0] * gridDim[1] * 1; try { this->InitExecNoInput(this->m_PixelCalculation, gridDim, m_BufferSize, sizeof(unsigned short)); } catch (const mitk::Exception& e) { MITK_ERROR << "Caught exception while initializing filter: " << e.what(); return; } if (m_Conf.DelayCalculationMethod == BeamformingSettings::DelayCalc::QuadApprox) m_DelayMultiplicatorRaw = pow(1 / (m_Conf.TimeSpacing*m_Conf.SpeedOfSound) * m_Conf.Pitch * (float)m_Conf.TransducerElements / (float)m_Conf.inputDim[0], 2) / 2; else if (m_Conf.DelayCalculationMethod == BeamformingSettings::DelayCalc::Spherical) m_DelayMultiplicatorRaw = 1 / (m_Conf.TimeSpacing*m_Conf.SpeedOfSound) * (m_Conf.Pitch*(float)m_Conf.TransducerElements); m_IsPAImage = m_Conf.isPhotoacousticImage; clErr = clSetKernelArg(this->m_PixelCalculation, 1, sizeof(cl_mem), &(this->m_UsedLines)); 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)); clErr |= clSetKernelArg(this->m_PixelCalculation, 6, sizeof(cl_char), &(this->m_IsPAImage)); clErr |= clSetKernelArg(this->m_PixelCalculation, 7, sizeof(cl_float), &(this->m_DelayMultiplicatorRaw)); CHECK_OCL_ERR(clErr); // execute the filter on a 3D NDRange if (!this->ExecuteKernelChunksInBatches(m_PixelCalculation, 2, m_ChunkSize, 16, 50)) mitkThrow() << "openCL Error when executing Kernel"; - // signalize the GPU-side data changed m_Output->Modified(GPU_DATA); } us::Module *mitk::OCLDelayCalculation::GetModule() { return us::GetModuleContext()->GetModule(); } bool mitk::OCLDelayCalculation::Initialize() { bool buildErr = true; cl_int clErr = 0; if (OclFilter::Initialize()) { if(m_Conf.DelayCalculationMethod == BeamformingSettings::DelayCalc::QuadApprox) this->m_PixelCalculation = clCreateKernel(this->m_ClProgram, "ckDelayCalculationQuad", &clErr); if (m_Conf.DelayCalculationMethod == BeamformingSettings::DelayCalc::Spherical) this->m_PixelCalculation = clCreateKernel(this->m_ClProgram, "ckDelayCalculationSphe", &clErr); buildErr |= CHECK_OCL_ERR(clErr); } return (OclFilter::IsInitialized() && buildErr); } \ No newline at end of file