kernelPBO.cu File Reference

#include <stdio.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <cutil_inline.h>
#include <cutil_math.h>

Classes
struct	float3x4
struct	Ray
struct	ViewParam

Macros
#define	PI   3.1415926535897932f

Typedefs
typedef unsigned char	VolumeType

Enumerations
enum	ColorMask { R_MASK = 0x000000FF, G_MASK = 0x0000FF00, B_MASK = 0x00FF0000, A_MASK = 0xFF000000 }

Functions
void	checkCUDAError (const char *msg)
void	init_cuda (void *h_volume, cudaExtent volumeSize, int dataset_nr)
	Implementation by kernelPBO.cu (compiled with nvcc)
void	freeCudaDofBuffers ()
void	init_Cuda_DoF (int image_width, int image_height)
void	freeCudaBuffers ()
__device__ float4	calculateGradientZucker (float x, float y, float z, float i)
__device__ int	intersectBox (Ray r, float3 boxmin, float3 boxmax, float *tnear, float *tfar)
__device__ float3	mul (const float3x4 &M, const float3 &v)
__device__ float4	mul (const float3x4 &M, const float4 &v)
__device__ uint	rgbaFloatToInt (float4 rgba)
__global__ void	kernel_volren (uint *d_output, uint imageW, uint imageH, float density, float brightness, float transferOffset, float transferScale, dim3 volumeSize)
__global__ void	kernel_composit (uint *d_output, float density, float brightness, uint imageW, uint imageH)
__global__ void	kernel_composit1 (uint *d_output, float brightness, bool FRONT_TO_BACK, uint imageW, uint imageH)
__global__ void	kernel_blur (float4 *d_nextBuffer, dim3 volumeSize, uint imageW, uint imageH, float density, float transferOffset, float transferScale, int i, float alpha, size_t pitch, const bool FRONT_TO_BACK, int fp, float d, float tstep)
void	copyInvViewMatrix (float *invViewMatrix, size_t sizeofMatrix)
void	swapBuffers (bool FRONT_TO_BACK, uint imageW, uint imageH)
__global__ void	kernel_calculateProxyGeometry (float4 *d_proxyStart, float4 *d_proxyEnd, size_t pitch, uint imageW, uint imageH, int *d_tnear, int *d_tfar)
void	calculateProxyGeometry (dim3 gridSize, dim3 blockSize, uint imageW, uint imageH, const float tstep, float4 focusPoint, float4 camPos)
void	launch_DepthOfField (dim3 gridSize, dim3 blockSize, dim3 volumeSize, uint *d_output, uint imageW, uint imageH, float density, float brightness, float transferOffset, float transferScale, float alpha, bool viewChanged, float4 focusPoint, float4 camPos, float tstep)
void	launch_volren (dim3 gridSize, dim3 blockSize, dim3 volumeSize, uint *d_output, uint imageW, uint imageH, float density, float brightness, float transferOffset, float transferScale)

Variables
cudaArray *	d_volumeArray = 0
cudaArray *	d_transferFuncArray
texture< VolumeType, 3, cudaReadModeNormalizedFloat >	tex
texture< float4, 1, cudaReadModeElementType >	transferTex
float4 *	d_frontBuffer
float4 *	d_backBuffer
float4 *	d_proxyStart
float4 *	d_proxyEnd
float4 *	d_nextBuffer
size_t	pitch_prev
size_t	pitch_next
bool *	d_finished_array
texture< float4, 2 >	tex_frontBuffer
texture< float4, 2 >	tex_backBuffer
texture< float4, 2 >	tex_proxyStart
texture< float4, 2 >	tex_proxyEnd
__constant__ float3x4	c_invViewMatrix
ViewParam	proxy

---

Macro Definition Documentation

#define PI   3.1415926535897932f

---

Typedef Documentation

typedef unsigned char VolumeType

---

Enumeration Type Documentation

enum ColorMask

Enumerator:

R_MASK
G_MASK
B_MASK
A_MASK

---

Function Documentation

__device__ float4 calculateGradientZucker	(	float	x,
		float	y,
		float	z,
		float	i
	)

calculate the Gradient as approximation for the Normal Vector

void calculateProxyGeometry	(	dim3	gridSize,
		dim3	blockSize,
		uint	imageW,
		uint	imageH,
		const float	tstep,
		float4	focusPoint,
		float4	camPos
	)

find out which of the slices is the focal plane (according to focusPoint (in world coordinates)

void checkCUDAError

(

const char *

msg

)

void copyInvViewMatrix	(	float *	invViewMatrix,
		size_t	sizeofMatrix
	)

void freeCudaBuffers

(

)

void freeCudaDofBuffers

(

)

free DoF Buffers and textures

void init_cuda	(	void *	h_volume,
		cudaExtent	volumeSize,
		int	dataset_nr
	)

Implementation by kernelPBO.cu (compiled with nvcc)

void init_Cuda_DoF	(	int	image_width,
		int	image_height
	)

Initialize DoF Buffers and textures

__device__ int intersectBox	(	Ray	r,
		float3	boxmin,
		float3	boxmax,
		float *	tnear,
		float *	tfar
	)

__global__ void kernel_blur	(	float4 *	d_nextBuffer,
		dim3	volumeSize,
		uint	imageW,
		uint	imageH,
		float	density,
		float	transferOffset,
		float	transferScale,
		int	i,
		float	alpha,
		size_t	pitch,
		const bool	FRONT_TO_BACK,
		int	fp,
		float	d,
		float	tstep
	)

kernel executed per slice for consecutive blurring for depth of field volume rendering by taking samples of previous rendered slices

__global__ void kernel_calculateProxyGeometry	(	float4 *	d_proxyStart,
		float4 *	d_proxyEnd,
		size_t	pitch,
		uint	imageW,
		uint	imageH,
		int *	d_tnear,
		int *	d_tfar
	)

__global__ void kernel_composit	(	uint *	d_output,
		float	density,
		float	brightness,
		uint	imageW,
		uint	imageH
	)

__global__ void kernel_composit1	(	uint *	d_output,
		float	brightness,
		bool	FRONT_TO_BACK,
		uint	imageW,
		uint	imageH
	)

__global__ void kernel_volren	(	uint *	d_output,
		uint	imageW,
		uint	imageH,
		float	density,
		float	brightness,
		float	transferOffset,
		float	transferScale,
		dim3	volumeSize
	)

void launch_DepthOfField	(	dim3	gridSize,
		dim3	blockSize,
		dim3	volumeSize,
		uint *	d_output,
		uint	imageW,
		uint	imageH,
		float	density,
		float	brightness,
		float	transferOffset,
		float	transferScale,
		float	alpha,
		bool	viewChanged,
		float4	focusPoint,
		float4	camPos,
		float	tstep
	)

launches the Depth of Field volume rendering cuda kernel

void launch_volren	(	dim3	gridSize,
		dim3	blockSize,
		dim3	volumeSize,
		uint *	d_output,
		uint	imageW,
		uint	imageH,
		float	density,
		float	brightness,
		float	transferOffset,
		float	transferScale
	)

launches the ordinary volume rendering cuda kernel

__device__ float3 mul	(	const float3x4 &	M,
		const float3 &	v
	)

__device__ float4 mul	(	const float3x4 &	M,
		const float4 &	v
	)

__device__ uint rgbaFloatToInt

(

float4

rgba

)

void swapBuffers	(	bool	FRONT_TO_BACK,
		uint	imageW,
		uint	imageH
	)

---

Variable Documentation

__constant__ float3x4 c_invViewMatrix

float4* d_backBuffer

bool* d_finished_array

float4* d_frontBuffer

float4* d_nextBuffer

float4* d_proxyEnd

float4* d_proxyStart

cudaArray* d_transferFuncArray

cudaArray* d_volumeArray = 0

size_t pitch_next

size_t pitch_prev

ViewParam proxy

texture<VolumeType, 3, cudaReadModeNormalizedFloat> tex

texture<float4, 2> tex_backBuffer

2D texture which stores the previous rendered slice buffer (back to front)

texture<float4, 2> tex_frontBuffer

2D texture which stores the previous rendered slice buffer (front to back)

texture<float4, 2> tex_proxyEnd

2D Texture storing the (Eyeray.d.x, Eyeray.d.y, Eyeray.d.z, tfar) for each pixel (updated every time the viewing parameters change)

texture<float4, 2> tex_proxyStart

2D Texture storing the (Eyeray.o.x, Eyeray.o.y, Eyeray.o.z, tnear) for each pixel (updated every time the viewing parameters change)

texture<float4, 1, cudaReadModeElementType> transferTex