Change Kernel not to keep internal copies of host values

src/c/Cuda/Kernel.cu

 #include "Kernel.cuh"
 #include "CudaUtilities.h"
 
-#ifndef CUDA_CONST_KERNEL
-#define CUDA_CONST_KERNEL
-__constant__ float cudaConstKernel[CONST_KERNEL_NUM_EL];
-#endif
+__constant__ float CUDA_KERNEL[CONST_KERNEL_NUM_EL];
 
-__host__ Kernel::Kernel(Vec<size_t> dimensions, float* values, int deviceIn, size_t startOffset/* = 0*/)
+__host__ Kernel::Kernel(Vec<size_t> dimensions, float* values, int curDevice, size_t startOffset/* = 0*/)
 {
-	load(dimensions, values, deviceIn, startOffset);
+	HANDLE_ERROR(cudaSetDevice(curDevice));
+	load(dimensions, values, startOffset);
 }
 
-__device__ Kernel::Kernel(const Kernel& other)
+__host__ __device__ Kernel::Kernel(Kernel& other)
 {
-	dims = other.dims;
-	kernel = other.kernel;
-	cleanUpHost = other.cleanUpHost;
-	cudaKernel = other.cudaKernel;
-	cleanUpDevice = other.cleanUpDevice;
+	init();
+
+	device = other.device;
+	dims = other.getDims();
+	cudaKernel = other.getDevicePtr();
+	cleanUpDevice = false;
 }
 
+__host__ Kernel::Kernel(ImageContainer<float> kernelIn, int curDevice)
+{
+	HANDLE_ERROR(cudaSetDevice(curDevice));
+	load(kernelIn.getSpatialDims(), kernelIn.getPtr());
+}
 
-__host__ Kernel::Kernel(ImageContainer<float> kernelIn, int deviceIn)
+__host__ Kernel::Kernel(size_t dimensions, float* values, int curDevice, size_t startOffset /*= 0*/)
 {
-	load(kernelIn.getSpatialDims(), kernelIn.getPtr(), deviceIn);
+	HANDLE_ERROR(cudaSetDevice(curDevice));
+	load(Vec<size_t>(dimensions, 1, 1), values, startOffset);
 }
 
+__device__ float Kernel::operator[](size_t idx)
+{
+	float val = cudaKernel[idx];
+	return val;
+}
 
-__host__ void Kernel::load(Vec<size_t> dimensions, float* values, int deviceIn, size_t startOffset/* = 0*/)
+__device__ float Kernel::operator()(Vec<size_t> coordinate)
+{
+	size_t idx = dims.linearAddressAt(coordinate);
+	float val = cudaKernel[idx];
+	return val;
+}
+
+__host__ Kernel& Kernel::operator=(const Kernel& other)
+{
+	dims = other.dims;
+	cudaKernel = other.cudaKernel;
+	cleanUpDevice = other.cleanUpDevice;
+
+	return *this;
+}
+
+__host__ void Kernel::load(Vec<size_t> dimensions, float* values, size_t startOffset/* = 0*/)
 {
 	init();
-	device = deviceIn;
-	cudaSetDevice(device);
 
 	dims = dimensions;
-
+	float* kernel = NULL;
+	bool cleanKernel = false;
 	if (values == NULL)
 	{
-		setOnes();
+		setOnes(&kernel);
+		cleanKernel = true;
 	}
 	else
 	{
@@ -46,9 +72,9 @@ __host__ void Kernel::load(Vec<size_t> dimensions, float* values, int deviceIn,
 
 	if (dimensions.product()+startOffset < CONST_KERNEL_NUM_EL)
 	{
-		HANDLE_ERROR(cudaGetSymbolAddress((void**)&cudaKernel, cudaConstKernel));
+		HANDLE_ERROR(cudaGetSymbolAddress((void**)&cudaKernel, CUDA_KERNEL));
 		cudaKernel += startOffset;
-		HANDLE_ERROR(cudaMemcpyToSymbol(cudaConstKernel, kernel, sizeof(float)*dims.product()));
+		HANDLE_ERROR(cudaMemcpyToSymbol(CUDA_KERNEL, kernel, sizeof(float)*dims.product()));
 	}
 	else
 	{
@@ -56,16 +82,13 @@ __host__ void Kernel::load(Vec<size_t> dimensions, float* values, int deviceIn,
 		HANDLE_ERROR(cudaMemcpy(cudaKernel, values, sizeof(float)*dims.product(),cudaMemcpyHostToDevice));
 		cleanUpDevice = true;
 	}
+
+	if (cleanKernel)
+		delete[] kernel;
 }
 
 __host__ void Kernel::clean()
 {
-	if (cleanUpHost)
-	{
-		delete[] kernel;
-		cleanUpHost = false;
-	}
-
 	if (cleanUpDevice)
 	{
 		cudaFree(cudaKernel);
@@ -75,36 +98,30 @@ __host__ void Kernel::clean()
 	init();
 }
 
-
-__device__ float Kernel::operator[](size_t idx)
-{
-	return cudaKernel[idx];
-}
-
-
-__device__ float Kernel::operator()(Vec<size_t> coordinate)
-{
-	return cudaKernel[dims.linearAddressAt(coordinate)];
-}
-
-
-__host__ void Kernel::init()
+__host__ __device__ void Kernel::init()
 {
 	dims = Vec<size_t>(0);
-	kernel = NULL;
-	cleanUpHost = false;
 	cudaKernel = NULL;
 	cleanUpDevice = false;
 }
 
-
-__host__ void Kernel::setOnes()
+__host__ void Kernel::setOnes(float** kernel)
 {
-	kernel = new float[dims.product()];
+	*kernel = new float[dims.product()];
 	for (int i = 0; i < dims.product(); ++i)
-		kernel[i] = 1.0f;
-
-	cleanUpHost = true;
+		(*kernel)[i] = 1.0f;
 }
 
+__host__ Kernel Kernel::getOffsetCopy(Vec<size_t> dimensions, size_t startOffset /*= 0*/)
+{
+	Kernel kernOut;
+	kernOut.init();
+
+	if (dims.product() < startOffset + dimensions.product())
+		std::runtime_error("Trying to make a Kernel that access outside of the original memory space!");
 
+	kernOut.dims = dimensions;
+	kernOut.cudaKernel = cudaKernel + startOffset;
+
+	return kernOut;
+}

src/c/Cuda/Kernel.cuh

@@ -8,30 +8,31 @@
 class Kernel
 {
 public:
-	__host__ Kernel(Vec<size_t> dimensions, float* values, int deviceIn, size_t startOffset=0);
-	__host__ Kernel(ImageContainer<float> kernelIn, int deviceIn);
-	__host__ __device__ Kernel(const Kernel& other);
+	__host__ Kernel(Vec<size_t> dimensions, float* values, int curDevice, size_t startOffset = 0);
+	__host__ Kernel(size_t dimensions, float* values, int curDevice, size_t startOffset = 0);
+	__host__ Kernel(ImageContainer<float> kernelIn, int curDevice);
+	__host__ __device__ Kernel(Kernel& other);
+	__host__ __device__ ~Kernel() { init(); }
 
-	__host__ __device__ ~Kernel() {}
+	__host__ Kernel& operator=(const Kernel& other);
 	__host__ void clean();
 
-	__host__ __device__ Vec<size_t> getDimensions() const { return dims; }
+	__host__ __device__ float* getDevicePtr() { return cudaKernel; }
+	__host__ Kernel getOffsetCopy(Vec<size_t> dimensions, size_t startOffset = 0);
+	__host__ __device__ Vec<size_t> getDims() const { return dims; }
 	__device__ float operator[](size_t idx);
 	__device__ float operator()(Vec<size_t> coordinate);
 	
 private:
-	__host__ __device__ Kernel();
+	__host__ __device__ Kernel() { init(); }
 
-	__host__ void load(Vec<size_t> dimensions, float* values, int deviceIn, size_t startOffset=0);
+	__host__ void load(Vec<size_t> dimensions, float* values, size_t startOffset=0);
 
-	__host__ void init();
-	__host__ void setOnes();
+	__host__ __device__ void init();
+	__host__ void setOnes(float** kernel);
 
 	Vec<size_t> dims;
 
-	float* kernel;
-	bool cleanUpHost;
-
 	float* cudaKernel;
 	bool cleanUpDevice;