Start of Kernel itterator class

src/c/Cuda/KernelIterator.cu

+#include "KernelIterator.cuh"
+
+#include <limits.h>
+#include <float.h>
+
+__device__ KernelIterator::KernelIterator(Vec<size_t> inputPos, Vec<size_t> inputSize, Vec<size_t> kernelSize)
+{
+	Vec<float> inputCoordinate(inputPos);
+
+	// If kernel size is odd, this will be integer, else we will get a real number
+	Vec<float> kernelHalfSize = (Vec<float>(kernelSize)-1)/2.0f;
+	
+	// This puts the first kernel position in the center of the voxel
+	// i.e. 1 is the second voxel position but 1.5 is an interpolated position between the first and second positions
+	inputStartCoordinate = inputCoordinate - kernelHalfSize;
+
+	// The negative here will make the kernel start w/in if it went out of bounds
+	// e.g. a -1 will start at the kernel's second position 
+	kernelStartIdx = Vec<size_t>::max(Vec<size_t>(0), Vec<size_t>(-inputStartCoordinate));
+
+	// This should be the last place in the image that will be considered
+	Vec<float> imageEndCoordinate = inputStartCoordinate + Vec<float>(kernelSize-1);
+	Vec<float> outOfBoundsAmount = imageEndCoordinate - Vec<float>(inputSize-1);
+
+	kernelEndIdx = kernelSize - Vec<size_t>(Vec<float>::max(1, outOfBoundsAmount +1));
+
+	iterator = Vec<size_t>(0);
+	isEnd = false;
+}
+
+__device__ __host__ KernelIterator::~KernelIterator()
+{
+	isEnd = true;
+	iterator = Vec<size_t>(ULLONG_MAX);
+	inputStartCoordinate = Vec<float>(FLT_MAX_EXP);
+	kernelStartIdx = Vec<size_t>(ULLONG_MAX);
+	kernelEndIdx = Vec<size_t>(0);
+}
+
+__device__ KernelIterator& KernelIterator::operator++()
+{
+	if(++iterator.x>kernelEndIdx.x)
+	{
+		iterator.x = kernelStartIdx.x;
+		
+		if(++iterator.y>kernelEndIdx.y)
+		{
+			iterator.y = kernelStartIdx.y;
+			if(++iterator.z>kernelEndIdx.z)
+			{
+				iterator.z = kernelEndIdx.z;
+				isEnd = true;
+			}
+		}
+	}
+
+	return *this;
+}
+
+__device__ Vec<float> KernelIterator::getImageCoordinate()
+{
+	return inputStartCoordinate + Vec<float>(iterator);
+}
+
+__device__ Vec<size_t> KernelIterator::getKernelIdx()
+{
+	return iterator;
+}
\ No newline at end of file

src/c/Cuda/KernelIterator.cuh

+#pragma once
+
+#include "Vec.h"
+
+class KernelIterator
+{
+public:
+	__device__ KernelIterator(Vec<size_t> inputPos, Vec<size_t> inputSize, Vec<size_t> kernelSize);
+	__device__ __host__ ~KernelIterator();
+
+	__device__ KernelIterator& operator++();
+	__device__ bool end() { return isEnd; }
+	__device__ Vec<float> getImageCoordinate();
+	__device__ Vec<size_t> getKernelIdx();
+	
+private:
+	__device__ __host__ KernelIterator() {}
+
+	// This is the first coordinate that to be used in the image
+	// If the coordinate is non-integer, then the image should be interpolated
+	Vec<float> inputStartCoordinate;
+
+	// This is the first index of the kernel to use
+	Vec<size_t> kernelStartIdx;
+
+	// This is the last index of the kernel to use (e.g. end+1 is out of bounds)
+	Vec<size_t> kernelEndIdx;
+
+	// This indicates the current position
+	Vec<size_t> iterator;
+
+	bool isEnd;
+};
+

src/c/CudaImageProcessor_dll.vcxproj

@@ -72,6 +72,7 @@
     <ClInclude Include="Cuda\CWrappers.h" />
     <ClInclude Include="Cuda\Defines.h" />
     <ClInclude Include="Cuda\ImageChunk.cuh" />
+    <ClInclude Include="Cuda\KernelIterator.cuh" />
     <ClInclude Include="Cuda\Vec.h" />
   </ItemGroup>
   <ItemGroup>
@@ -84,6 +85,7 @@
     <CudaCompile Include="Cuda\CudaUtilities.cu" />
     <CudaCompile Include="Cuda\CWrappers.cu" />
     <CudaCompile Include="Cuda\ImageChunk.cu" />
+    <CudaCompile Include="Cuda\KernelIterator.cu" />
   </ItemGroup>
   <PropertyGroup Label="Globals">
     <ProjectGuid>{F204D99B-C454-420D-BD1B-5FE4D9D07078}</ProjectGuid>
@@ -187,6 +189,7 @@ copy "$(CudaToolkitBinDir)\cudart*.dll" "$(OutDir)"</Command>
       <PtxAsOptionV>false</PtxAsOptionV>
       <MaxRegCount>
       </MaxRegCount>
+      <GenerateRelocatableDeviceCode>true</GenerateRelocatableDeviceCode>
     </CudaCompile>
     <Lib>
       <AdditionalDependencies>cudart_static.lib</AdditionalDependencies>
@@ -239,6 +242,7 @@ copy "$(CudaToolkitBinDir)\cudart*.dll" "$(OutDir)"</Command>
       <CodeGeneration>compute_30,sm_30;compute_35,sm_35;compute_50,sm_50;</CodeGeneration>
       <MaxRegCount>
       </MaxRegCount>
+      <GenerateRelocatableDeviceCode>true</GenerateRelocatableDeviceCode>
     </CudaCompile>
     <Lib>
       <AdditionalDependencies>cudart_static.lib</AdditionalDependencies>

src/c/CudaImageProcessor_dll.vcxproj.filters

@@ -146,6 +146,9 @@
     <ClInclude Include="Cuda\CudaLoG.cuh">
       <Filter>Header Files</Filter>
     </ClInclude>
+    <ClInclude Include="Cuda\KernelIterator.cuh">
+      <Filter>Header Files</Filter>
+    </ClInclude>
   </ItemGroup>
   <ItemGroup>
     <ClCompile Include="Cuda\CHelpers.cpp">
@@ -165,5 +168,8 @@
     <CudaCompile Include="Cuda\ImageChunk.cu">
       <Filter>Source Files</Filter>
     </CudaCompile>
+    <CudaCompile Include="Cuda\KernelIterator.cu">
+      <Filter>Source Files</Filter>
+    </CudaCompile>
   </ItemGroup>
 </Project>
\ No newline at end of file