Added identity copy filter for testing purposes

Identity filter copies data to GPU then back to output image.

src/c/Cuda/CWrappers.cu

@@ -9,6 +9,7 @@
 #include "CudaGaussian.cuh"
 #include "CudaGetMinMax.cuh"
 #include "CudaHighPassFilter.cuh"
+#include "CudaIdentityFilter.cuh"
 #include "CudaLoG.cuh"
 #include "CudaMaxFilter.cuh"
 #include "CudaMedianFilter.cuh"
@@ -372,6 +373,52 @@ void gaussian(const ImageContainer<bool> imageIn, ImageContainer<bool>& imageOut
  }
 
 
+ void identityFilter(const ImageContainer<bool> imageIn, ImageContainer<bool>& imageOut, int device /*= -1*/)
+ {
+	 cIdentityFilter(imageIn, imageOut, device);
+ }
+
+ void identityFilter(const ImageContainer<char> imageIn, ImageContainer<char>& imageOut, int device /*= -1*/)
+ {
+	 cIdentityFilter(imageIn, imageOut, device);
+ }
+
+ void identityFilter(const ImageContainer<unsigned char> imageIn, ImageContainer<unsigned char>& imageOut, int device /*= -1*/)
+ {
+	 cIdentityFilter(imageIn, imageOut, device);
+ }
+
+ void identityFilter(const ImageContainer<short> imageIn, ImageContainer<short>& imageOut, int device /*= -1*/)
+ {
+	 cIdentityFilter(imageIn, imageOut, device);
+ }
+
+ void identityFilter(const ImageContainer<unsigned short> imageIn, ImageContainer<unsigned short>& imageOut, int device /*= -1*/)
+ {
+	 cIdentityFilter(imageIn, imageOut, device);
+ }
+
+ void identityFilter(const ImageContainer<int> imageIn, ImageContainer<int>& imageOut, int device /*= -1*/)
+ {
+	 cIdentityFilter(imageIn, imageOut, device);
+ }
+
+ void identityFilter(const ImageContainer<unsigned int> imageIn, ImageContainer<unsigned int>& imageOut, int device /*= -1*/)
+ {
+	 cIdentityFilter(imageIn, imageOut, device);
+ }
+
+ void identityFilter(const ImageContainer<float> imageIn, ImageContainer<float>& imageOut, int device /*= -1*/)
+ {
+	 cIdentityFilter(imageIn, imageOut, device);
+ }
+
+ void identityFilter(const ImageContainer<double> imageIn, ImageContainer<double>& imageOut, int device /*= -1*/)
+ {
+	 cIdentityFilter(imageIn, imageOut, device);
+ }
+
+
  void LoG(const ImageContainer<bool> imageIn, ImageContainer<float>& imageOut, Vec<double> sigmas, int device /*= -1*/)
  {
 	 cLoG(imageIn, imageOut, sigmas, device);

src/c/Cuda/CWrappers.h

@@ -93,6 +93,16 @@ IMAGE_PROCESSOR_API void highPassFilter(const ImageContainer<unsigned int> image
 IMAGE_PROCESSOR_API void highPassFilter(const ImageContainer<float> imageIn, ImageContainer<float>& imageOut, Vec<double> sigmas, int device = -1);
 IMAGE_PROCESSOR_API void highPassFilter(const ImageContainer<double> imageIn, ImageContainer<double>& imageOut, Vec<double> sigmas, int device = -1);
 
+IMAGE_PROCESSOR_API void identityFilter(const ImageContainer<bool> imageIn, ImageContainer<bool>& imageOut, int device = -1);
+IMAGE_PROCESSOR_API void identityFilter(const ImageContainer<char> imageIn, ImageContainer<char>& imageOut, int device = -1);
+IMAGE_PROCESSOR_API void identityFilter(const ImageContainer<unsigned char> imageIn, ImageContainer<unsigned char>& imageOut, int device = -1);
+IMAGE_PROCESSOR_API void identityFilter(const ImageContainer<short> imageIn, ImageContainer<short>& imageOut, int device = -1);
+IMAGE_PROCESSOR_API void identityFilter(const ImageContainer<unsigned short> imageIn, ImageContainer<unsigned short>& imageOut, int device = -1);
+IMAGE_PROCESSOR_API void identityFilter(const ImageContainer<int> imageIn, ImageContainer<int>& imageOut, int device = -1);
+IMAGE_PROCESSOR_API void identityFilter(const ImageContainer<unsigned int> imageIn, ImageContainer<unsigned int>& imageOut, int device = -1);
+IMAGE_PROCESSOR_API void identityFilter(const ImageContainer<float> imageIn, ImageContainer<float>& imageOut, int device = -1);
+IMAGE_PROCESSOR_API void identityFilter(const ImageContainer<double> imageIn, ImageContainer<double>& imageOut, int device = -1);
+
 IMAGE_PROCESSOR_API void LoG(const ImageContainer<bool> imageIn, ImageContainer<float>& imageOut, Vec<double> sigmas, int device = -1);
 IMAGE_PROCESSOR_API void LoG(const ImageContainer<char> imageIn, ImageContainer<float>& imageOut, Vec<double> sigmas, int device = -1);
 IMAGE_PROCESSOR_API void LoG(const ImageContainer<unsigned char> imageIn, ImageContainer<float>& imageOut, Vec<double> sigmas, int device = -1);

src/c/Cuda/CudaIdentityFilter.cuh

+#pragma once
+
+#include "CudaImageContainer.cuh"
+#include "CudaDeviceImages.cuh"
+#include "CudaUtilities.h"
+#include "CudaDeviceInfo.h"
+#include "Kernel.cuh"
+#include "KernelIterator.cuh"
+#include "ImageDimensions.cuh"
+#include "ImageChunk.h"
+#include "Defines.h"
+#include "Vec.h"
+
+#include <cuda_runtime.h>
+#include <limits>
+#include <omp.h>
+
+template <class PixelTypeIn, class PixelTypeOut>
+__global__ void cudaIdentity(CudaImageContainer<PixelTypeIn> imageIn, CudaImageContainer<PixelTypeOut> imageOut)
+{
+	Vec<std::size_t> threadCoordinate;
+	GetThreadBlockCoordinate(threadCoordinate);
+
+	if ( threadCoordinate < imageIn.getDims() )
+	{
+		imageOut(threadCoordinate) = (PixelTypeOut)imageIn(threadCoordinate);
+	}
+}
+
+
+template <class PixelTypeIn, class PixelTypeOut>
+void cIdentityFilter(ImageContainer<PixelTypeIn> imageIn, ImageContainer<PixelTypeOut>& imageOut, int device = -1)
+{
+	const int NUM_BUFF_NEEDED = 2;
+
+	setUpOutIm<PixelTypeOut>(imageIn.getDims(), imageOut);
+
+	CudaDevices cudaDevs(cudaIdentity<PixelTypeIn, PixelTypeOut>, device);
+
+	Vec<std::size_t> kernDims(5*75+1);
+
+	std::size_t maxTypeSize = MAX(sizeof(PixelTypeIn), sizeof(PixelTypeOut));
+	std::vector<ImageChunk> chunks = calculateBuffers(imageIn.getDims(), NUM_BUFF_NEEDED, cudaDevs, maxTypeSize, kernDims);
+
+	Vec<std::size_t> maxDeviceDims;
+	setMaxDeviceDims(chunks, maxDeviceDims);
+
+	omp_set_num_threads(MIN(chunks.size(), cudaDevs.getNumDevices()));
+	#pragma omp parallel default(shared)
+	{
+		const int CUDA_IDX = omp_get_thread_num();
+		const int N_THREADS = omp_get_num_threads();
+		const int CUR_DEVICE = cudaDevs.getDeviceIdx(CUDA_IDX);
+
+		CudaDeviceImages<PixelTypeOut> deviceImages(NUM_BUFF_NEEDED, maxDeviceDims, CUR_DEVICE);
+
+		for ( int i = CUDA_IDX; i < chunks.size(); i += N_THREADS )
+		{
+			if ( !chunks[i].sendROI(imageIn, deviceImages.getCurBuffer()) )
+				std::runtime_error("Error sending ROI to device!");
+
+			deviceImages.setAllDims(chunks[i].getFullChunkSize());
+
+			cudaIdentity<<<chunks[i].blocks, chunks[i].threads>>>(*(deviceImages.getCurBuffer()), *(deviceImages.getNextBuffer()));
+			deviceImages.incrementBuffer();
+
+			chunks[i].retriveROI(imageOut, deviceImages.getCurBuffer());
+		}
+	}
+}

src/c/CudaImageProcessor.vcxproj

@@ -127,6 +127,7 @@
     <ClInclude Include="Cuda\CudaClosure.cuh" />
     <ClInclude Include="Cuda\CudaEntropyFilter.cuh" />
     <ClInclude Include="Cuda\CudaGetMinMax.cuh" />
+    <ClInclude Include="Cuda\CudaIdentityFilter.cuh" />
     <ClInclude Include="Cuda\CudaLoG.cuh" />
     <ClInclude Include="Cuda\CudaMeanAndVariance.cuh" />
     <ClInclude Include="Cuda\CudaMeanFilter.cuh" />

src/c/CudaImageProcessor.vcxproj.filters

@@ -155,6 +155,9 @@
     <ClInclude Include="Cuda\CudaVarFilter.cuh">
       <Filter>Header Files</Filter>
     </ClInclude>
+    <ClInclude Include="Cuda\CudaIdentityFilter.cuh">
+      <Filter>Header Files</Filter>
+    </ClInclude>
   </ItemGroup>
   <ItemGroup>
     <None Include="Cuda\ImageDimensions.cuh">

src/c/CudaMex.vcxproj

@@ -107,6 +107,7 @@ copy $(OutDir)CudaMex.dll "$(ProjectDir)Mex.mexw64"</Command>
     <ClCompile Include="Mex\MexElementWiseDifference.cpp" />
     <ClCompile Include="Mex\MexGetMinMax.cpp" />
     <ClCompile Include="Mex\MexEntropyFilter.cpp" />
+    <ClCompile Include="Mex\MexIdentityFilter.cpp" />
     <ClCompile Include="Mex\MexMedianFilter.cpp" />
     <ClCompile Include="Mex\MexClosure.cpp" />
     <ClCompile Include="Mex\MexCommand.cpp" />

src/c/CudaMex.vcxproj.filters

@@ -110,6 +110,9 @@
     <ClCompile Include="Mex\MexVarFilter.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
+    <ClCompile Include="Mex\MexIdentityFilter.cpp">
+      <Filter>Source Files</Filter>
+    </ClCompile>
   </ItemGroup>
   <ItemGroup>
     <None Include="Mex\CudaMex.def">

src/c/CudaPy3DLL.vcxproj

@@ -30,6 +30,7 @@
     <ClCompile Include="Python\PyWrapGaussian.cpp" />
     <ClCompile Include="Python\PyWrapGetMinMax.cpp" />
     <ClCompile Include="Python\PyWrapHighPassFilter.cpp" />
+    <ClCompile Include="Python\PyWrapIdentityFilter.cpp" />
     <ClCompile Include="Python\PyWrapLoG.cpp" />
     <ClCompile Include="Python\PyWrapMaxFilter.cpp" />
     <ClCompile Include="Python\PyWrapMeanFilter.cpp" />

src/c/CudaPy3DLL.vcxproj.filters

@@ -84,6 +84,9 @@
     <ClCompile Include="Python\PyWrapMultiplySum.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
+    <ClCompile Include="Python\PyWrapIdentityFilter.cpp">
+      <Filter>Source Files</Filter>
+    </ClCompile>
   </ItemGroup>
   <ItemGroup>
     <ClInclude Include="WrapCmds\CommandList.h">

src/c/Mex/MexIdentityFilter.cpp

+#include "MexCommand.h"
+#include "../Cuda/Vec.h"
+#include "../Cuda/CWrappers.h"
+#include "../Cuda/ImageDimensions.cuh"
+#include "../Cuda/ImageContainer.h"
+#include "MexKernel.h"
+
+void MexIdentityFilter::execute(int nlhs, mxArray* plhs[], int nrhs, const mxArray* prhs[]) const
+{
+	int device = -1;
+
+	ImageDimensions imageDims;
+	if ( mxIsLogical(prhs[0]) )
+	{
+		bool* imageInPtr, *imageOutPtr;
+		Script::setupImagePointers(prhs[0], &imageInPtr, imageDims, &plhs[0], &imageOutPtr);
+
+		ImageContainer<bool> imageIn(imageInPtr, imageDims);
+		ImageContainer<bool> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+
+	}
+	else if ( mxIsUint8(prhs[0]) )
+	{
+		unsigned char* imageInPtr, *imageOutPtr;
+		Script::setupImagePointers(prhs[0], &imageInPtr, imageDims, &plhs[0], &imageOutPtr);
+
+		ImageContainer<unsigned char> imageIn(imageInPtr, imageDims);
+		ImageContainer<unsigned char> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else if ( mxIsUint16(prhs[0]) )
+	{
+		unsigned short* imageInPtr, *imageOutPtr;
+		Script::setupImagePointers(prhs[0], &imageInPtr, imageDims, &plhs[0], &imageOutPtr);
+
+		ImageContainer<unsigned short> imageIn(imageInPtr, imageDims);
+		ImageContainer<unsigned short> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else if ( mxIsInt16(prhs[0]) )
+	{
+		short* imageInPtr, *imageOutPtr;
+		Script::setupImagePointers(prhs[0], &imageInPtr, imageDims, &plhs[0], &imageOutPtr);
+
+		ImageContainer<short> imageIn(imageInPtr, imageDims);
+		ImageContainer<short> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else if ( mxIsUint32(prhs[0]) )
+	{
+		unsigned int* imageInPtr, *imageOutPtr;
+		Script::setupImagePointers(prhs[0], &imageInPtr, imageDims, &plhs[0], &imageOutPtr);
+
+		ImageContainer<unsigned int> imageIn(imageInPtr, imageDims);
+		ImageContainer<unsigned int> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else if ( mxIsInt32(prhs[0]) )
+	{
+		int* imageInPtr, *imageOutPtr;
+		Script::setupImagePointers(prhs[0], &imageInPtr, imageDims, &plhs[0], &imageOutPtr);
+
+		ImageContainer<int> imageIn(imageInPtr, imageDims);
+		ImageContainer<int> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else if ( mxIsSingle(prhs[0]) )
+	{
+		float* imageInPtr, *imageOutPtr;
+		Script::setupImagePointers(prhs[0], &imageInPtr, imageDims, &plhs[0], &imageOutPtr);
+
+		ImageContainer<float> imageIn(imageInPtr, imageDims);
+		ImageContainer<float> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else if ( mxIsDouble(prhs[0]) )
+	{
+		double* imageInPtr, *imageOutPtr;
+		Script::setupImagePointers(prhs[0], &imageInPtr, imageDims, &plhs[0], &imageOutPtr);
+
+		ImageContainer<double> imageIn(imageInPtr, imageDims);
+		ImageContainer<double> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else
+	{
+		mexErrMsgTxt("Image type not supported!");
+	}
+}
+
+std::string MexIdentityFilter::check(int nlhs, mxArray* plhs[], int nrhs, const mxArray* prhs[]) const
+{
+	if ( nrhs != 2 )
+		return "Incorrect number of inputs!";
+
+	if ( nlhs != 1 )
+		return "Requires one output!";
+
+	std::size_t imgNumDims = mxGetNumberOfDimensions(prhs[0]);
+	if ( imgNumDims > 5 )
+		return "Image can have a maximum of five dimensions!";
+
+
+	return "";
+}
+
+void MexIdentityFilter::usage(std::vector<std::string>& outArgs, std::vector<std::string>& inArgs) const
+{
+	inArgs.push_back("arrayIn");
+	inArgs.push_back("[device]");
+	outArgs.push_back("arrayOut");
+}
+
+void MexIdentityFilter::help(std::vector<std::string>& helpLines) const
+{
+	helpLines.push_back("Identity Filter for testing. Copies image data to GPU memory and back into output image");
+
+	helpLines.push_back("\timageIn = This is a one to five dimensional array. The first three dimensions are treated as spatial.");
+	helpLines.push_back("\t\tThe spatial dimensions will have the kernel applied. The last two dimensions will determine");
+	helpLines.push_back("\t\thow to stride or jump to the next spatial block.");
+	helpLines.push_back("");
+
+	helpLines.push_back("\tdevice (optional) = Use this if you have multiple devices and want to select one explicitly.");
+	helpLines.push_back("\t\tSetting this to [] allows the algorithm to either pick the best device and/or will try to split");
+	helpLines.push_back("\t\tthe data across multiple devices.");
+	helpLines.push_back("");
+
+	helpLines.push_back("\timageOut = This will be an array of the same type and shape as the input array.");
+}

src/c/Python/PyWrapIdentityFilter.cpp

+#include "PyWrapCommand.h"
+
+#include "../Cuda/Vec.h"
+#include "../Cuda/CWrappers.h"
+#include "../Cuda/ImageDimensions.cuh"
+#include "../Cuda/ImageContainer.h"
+
+#include "PyKernel.h"
+
+
+
+const char PyWrapIdentityFilter::docString[] = "imageOut = HIP.IdentityFilter(imageIn,[device])\n\n"\
+"Identity Filter for testing. Copies image data to GPU memory and back into output image.\n"\
+"\timageIn = This is a one to five dimensional array. The first three dimensions are treated as spatial.\n"\
+"\t\tThe spatial dimensions will have the kernel applied. The last two dimensions will determine\n"\
+"\t\thow to stride or jump to the next spatial block.\n"\
+"\n"\
+"\tdevice (optional) = Use this if you have multiple devices and want to select one explicitly.\n"\
+"\t\tSetting this to [] allows the algorithm to either pick the best device and/or will try to split\n"\
+"\t\tthe data across multiple devices.\n"\
+"\n"\
+"\timageOut = This will be an array of the same type and shape as the input array.\n";
+
+
+PyObject* PyWrapIdentityFilter::execute(PyObject* self, PyObject* args)
+{
+	int device = -1;
+
+	PyObject* imIn;
+
+	if ( !PyArg_ParseTuple(args, "O!|i", &PyArray_Type, &imIn, &device) )
+		return nullptr;
+
+	if ( imIn == nullptr ) return nullptr;
+
+	PyArrayObject* imContig = (PyArrayObject*)PyArray_FROM_OTF(imIn, NPY_NOTYPE, NPY_ARRAY_IN_ARRAY);
+
+	ImageDimensions imageDims;
+	PyArrayObject* imOut = nullptr;
+
+	if ( PyArray_TYPE(imContig) == NPY_BOOL )
+	{
+		bool* imageInPtr, *imageOutPtr;
+
+		Script::setupImagePointers(imContig, &imageInPtr, imageDims, &imOut, &imageOutPtr);
+
+		ImageContainer<bool> imageIn(imageInPtr, imageDims);
+		ImageContainer<bool> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+
+	}
+	else if ( PyArray_TYPE(imContig) == NPY_UINT8 )
+	{
+		unsigned char* imageInPtr, *imageOutPtr;
+
+		Script::setupImagePointers(imContig, &imageInPtr, imageDims, &imOut, &imageOutPtr);
+
+		ImageContainer<unsigned char> imageIn(imageInPtr, imageDims);
+		ImageContainer<unsigned char> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else if ( PyArray_TYPE(imContig) == NPY_UINT16 )
+	{
+		unsigned short* imageInPtr, *imageOutPtr;
+
+		Script::setupImagePointers(imContig, &imageInPtr, imageDims, &imOut, &imageOutPtr);
+
+		ImageContainer<unsigned short> imageIn(imageInPtr, imageDims);
+		ImageContainer<unsigned short> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else if ( PyArray_TYPE(imContig) == NPY_INT16 )
+	{
+		short* imageInPtr, *imageOutPtr;
+
+		Script::setupImagePointers(imContig, &imageInPtr, imageDims, &imOut, &imageOutPtr);
+
+		ImageContainer<short> imageIn(imageInPtr, imageDims);
+		ImageContainer<short> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else if ( PyArray_TYPE(imContig) == NPY_UINT32 )
+	{
+		unsigned int* imageInPtr, *imageOutPtr;
+
+		Script::setupImagePointers(imContig, &imageInPtr, imageDims, &imOut, &imageOutPtr);
+
+		ImageContainer<unsigned int> imageIn(imageInPtr, imageDims);
+		ImageContainer<unsigned int> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else if ( PyArray_TYPE(imContig) == NPY_INT32 )
+	{
+		int* imageInPtr, *imageOutPtr;
+
+		Script::setupImagePointers(imContig, &imageInPtr, imageDims, &imOut, &imageOutPtr);
+
+		ImageContainer<int> imageIn(imageInPtr, imageDims);
+		ImageContainer<int> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else if ( PyArray_TYPE(imContig) == NPY_FLOAT )
+	{
+		float* imageInPtr, *imageOutPtr;
+
+		Script::setupImagePointers(imContig, &imageInPtr, imageDims, &imOut, &imageOutPtr);
+
+		ImageContainer<float> imageIn(imageInPtr, imageDims);
+		ImageContainer<float> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else if ( PyArray_TYPE(imContig) == NPY_DOUBLE )
+	{
+		double* imageInPtr, *imageOutPtr;
+
+		Script::setupImagePointers(imContig, &imageInPtr, imageDims, &imOut, &imageOutPtr);
+
+		ImageContainer<double> imageIn(imageInPtr, imageDims);
+		ImageContainer<double> imageOut(imageOutPtr, imageDims);
+
+		identityFilter(imageIn, imageOut, device);
+	}
+	else
+	{
+		PyErr_SetString(PyExc_RuntimeError, "Image type not supported.");
+
+		Py_XDECREF(imContig);
+		Py_XDECREF(imOut);
+
+		return nullptr;
+	}
+
+	Py_XDECREF(imContig);
+
+	return ((PyObject*)imOut);
+}

src/c/WrapCmds/CommandList.h

@@ -12,6 +12,7 @@ BEGIN_WRAP_COMMANDS
 	DEF_WRAP_COMMAND(Gaussian)
 	DEF_WRAP_COMMAND(GetMinMax)
 	DEF_WRAP_COMMAND(HighPassFilter)
+	DEF_WRAP_COMMAND(IdentityFilter)
 	DEF_WRAP_COMMAND(LoG)
 	DEF_WRAP_COMMAND(MaxFilter)
 	DEF_WRAP_COMMAND(MeanFilter)