ProcessingUtilsCuda.cu File Reference

#include "ProcessingUtilsCuda.h"
#include <cuda_runtime.h>
#include <vector_types.h>
#include <stdio.h>

Macros
#define	CHECK_CUDA(X)

Functions
__global__ void	convolve (float *resultBuffer, int width, int height, int depth, float *kern, int kW, int kH, int kD, int offset)
__global__ void	gradient (float4 *resultBuffer, int width, int height, int depth, int offset)
__device__ float	computeAlpha (float fx, float fy, float fz, float fxx, float fxy, float fxz, float fyx, float fyy, float fyz, float fzx, float fzy, float fzz)
__global__ void	curvature (float4 *resultBuffer, int width, int height, int depth, int offset, int offsetImage, int maxIdx)
int	cudaCurvature3D (float4 *gradient, float4 *resultBuffer, int width, int height, int depth)
int	cudaGradient3D (float *image, float4 *resultBuffer, int width, int height, int depth)
int	cudaConvolve3D (float *image, float *resultBuffer, int width, int height, int depth, float *kernel, int kW, int kH, int kD)

Variables
texture< float, 3, cudaReadModeElementType >	image_ref
texture< float4, 3, cudaReadModeElementType >	gradient_ref

---

Macro Definition Documentation

#define CHECK_CUDA

(

X

)

Value:

{                                               \
                cudaThreadSynchronize(); \
                cudaError_t retVal = (X);       \
                if (retVal != 0) {      \
                        fprintf(stderr, "CUDA Error: %d - %s on Line %d in File %s\n", retVal, cudaGetErrorString(retVal), __LINE__, __FILE__); \
                        return -1; \
                }       \
        }

---

Function Documentation

__device__ float computeAlpha	(	float	fx,
		float	fy,
		float	fz,
		float	fxx,
		float	fxy,
		float	fxz,
		float	fyx,
		float	fyy,
		float	fyz,
		float	fzx,
		float	fzy,
		float	fzz
	)

__global__ void convolve	(	float *	resultBuffer,
		int	width,
		int	height,
		int	depth,
		float *	kern,
		int	kW,
		int	kH,
		int	kD,
		int	offset
	)

  Convolves the volume with a kernel

Parameters:

resultBuffer	is the result buffer, in which the data should be written
width	is the width of the volume
height	is the height of the volume
depth	is the depth of the volume
kern	is the kernel
kW	is the kernel width
kH	is the kernel height
kD	is the kernel depth
offset	is an additional offset which is added to the index. This must be used, since we are convolving in XY-plane chunks along the z-axis for memory/block-size reasons

!!

int cudaConvolve3D	(	float *	image,
		float *	resultBuffer,
		int	width,
		int	height,
		int	depth,
		float *	kernel,
		int	kW,
		int	kH,
		int	kD
	)

  Convolves the 3D volume with cuda

Parameters:

image	is the 3D volume
resultBuffer	is a buffer for the result
width	is the width of the volume
height	is the height of the volume
depth	is the depth of the volume
kernel	is the convolution-kernel
kW	is the kernel width
kH	is the kernel height
kD	is the kernel depth

int cudaCurvature3D	(	float4 *	gradient,
		float4 *	resultBuffer,
		int	width,
		int	height,
		int	depth
	)

  Calculates the curvature with cuda

Parameters:

gradient	is the gradient image
resultBuffer	is the result buffer in which the result should be written
width	is the width of the volume
height	is the height of the volume
depth	is the depth of the volume

int cudaGradient3D	(	float *	image,
		float4 *	resultBuffer,
		int	width,
		int	height,
		int	depth
	)

  Calculates the gradient with cuda

Parameters:

image	is the volume-'image' whose gradient should be calculated
resultBuffer	is a buffer in which the gradient will be stored
width	is the width of the volume
height	is the height of the volume
depth	is the depth of the volume

__global__ void curvature	(	float4 *	resultBuffer,
		int	width,
		int	height,
		int	depth,
		int	offset,
		int	offsetImage,
		int	maxIdx
	)

  Calculates the curvature with cuda. The result image is only bound in XY-slices! So we must subtract offset from it!

Furthermore the gradient image is bound in chunks. The chunks start with offsetImage and end with maxId.

Parameters:

resultBuffer	is the buffer in which the results are written.
width	is the width of the buffer
height	is the height of the buffer
depth	is the depth of the buffer
offset	is an additional offset, since we are processing in XY-slices
offsetImage	is an offset in the gradient image
maxIdx	is basically the limit to which the gradient image is bound. Blocks executing after this image should not be run (or they accessing unbound gradient-chunks)

__global__ void gradient	(	float4 *	resultBuffer,
		int	width,
		int	height,
		int	depth,
		int	offset
	)

  Calculates the gradient.

Parameters:

resultBuffer	is the result image buffer
width	is the volume width
height	is the volume height
depth	is the volume depth
offset	is an additional index-offset since we are convolving the gradient in XY-plane-chunks for memory/block-size reasons

---

Variable Documentation

texture<float4, 3, cudaReadModeElementType> gradient_ref

texture<float, 3, cudaReadModeElementType> image_ref