<commit_before>/*===================================================================

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 <mitkLogMacros.h>
#include <mitkConfig.h>

//usService
#include "usServiceReference.h"
#include <usServiceRegistration.h>
#include <usModuleContext.h>
#include <usGetModuleContext.h>
#include <usModule.h>
#include <usModuleResource.h>
#include <usModuleResourceStream.h>

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<OclResourceService> ref = GetModuleContext()->GetServiceReference<OclResourceService>();
    OclResourceService* resources = GetModuleContext()->GetService<OclResourceService>(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::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::Initialize()
{
  us::ServiceReference<OclResourceService> ref = GetModuleContext()->GetServiceReference<OclResourceService>();
  OclResourceService* resources = GetModuleContext()->GetService<OclResourceService>(ref);<commit_msg>the chunked openCL execution waits for 16 chunks to finish until it queues the next ones, so that the system gets some time to check whether the GPU still responds<commit_after>/*===================================================================

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 <mitkLogMacros.h>
#include <mitkConfig.h>

//usService
#include "usServiceReference.h"
#include <usServiceRegistration.h>
#include <usModuleContext.h>
#include <usGetModuleContext.h>
#include <usModule.h>
#include <usModuleResource.h>
#include <usModuleResourceStream.h>

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<OclResourceService> ref = GetModuleContext()->GetServiceReference<OclResourceService>();
    OclResourceService* resources = GetModuleContext()->GetService<OclResourceService>(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::ExecuteKernelChunks( cl_kernel kernel, unsigned int workSizeDim, size_t* chunksDim )
{
  size_t offset[3] ={0, 0, 0};
  cl_int clErr = 0;

  unsigned int currentChunk = 0;
  cl_event* waitFor = new cl_event[16];

  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 % 16 == 0 && currentChunk != 0)
        {
          clWaitForEvents(16, &waitFor[0]);
          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%16]);
        }
        currentChunk++;
        //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::Initialize()
{
  us::ServiceReference<OclResourceService> ref = GetModuleContext()->GetServiceReference<OclResourceService>();
  OclResourceService* resources = GetModuleContext()->GetService<OclResourceService>(ref);
{'fraction_non_alphanumeric': 0.15706670133922768, 'fraction_numerical': 0.012872188272006242, 'mean_word_length': 4.027450980392157, 'pattern_counts': {'":': 0, '<': 45, '<?xml version=': 0, '>': 51, 'https://': 0, 'lorem ipsum': 0, 'www.': 2, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 7, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '6696574', 'n_tokens_mistral': 2527, 'n_tokens_neox': 2459, 'n_words': 511}