Main Page   Namespace List   Class Hierarchy   Alphabetical List   Compound List   File List   Namespace Members   Compound Members   File Members   Related Pages  

CUDAWatershed.cu File Reference

CUDA-accelerated Watershed image segmentation. More...

#include <cstdio>
#include "CUDAWatershed.h"
#include "ProfileHooks.h"

Go to the source code of this file.

Defines

#define WATERSHED_INTERNAL   1
#define px_offset_d   (1)
#define py_offset_d   (width)
#define pz_offset_d   (heightWidth)
#define nx_offset_d   (-1)
#define ny_offset_d   (-width)
#define nz_offset_d   (-heightWidth)
#define nx_ny_offset_d   (-1 - width)
#define nx_py_offset_d   (-1 + width)
#define px_py_offset_d   (1 + width)
#define px_ny_offset_d   (1 - width)
#define px_pz_offset_d   (1 + heightWidth)
#define nx_nz_offset_d   (-1 - heightWidth)
#define nx_pz_offset_d   (-1 + heightWidth)
#define px_nz_offset_d   (1 - heightWidth)
#define py_pz_offset_d   (width + heightWidth)
#define ny_nz_offset_d   (-width - heightWidth)
#define ny_pz_offset_d   (-width + heightWidth)
#define py_nz_offset_d   (width - heightWidth)
#define GPU_WARP_UPDATE
#define BLOCK_X_DIM   64
#define BLOCK_Y_DIM   2
#define BLOCK_Z_DIM   2
#define INST_DESTROY_GPU(G_T)
#define INST_INIT_GPU_ON_DEV(G_T)
#define INST_INIT_GPU(G_T)
#define INST_UPDATE_CUDA(G_T)
#define DIV_BY_32   5
#define GIGABYTE   (1024.0f*1024.0f*1024.0f)
#define CALCULATE_NEIGHBORS_CUDA(offset_str)
#define UPDATE_VOXEL_CUDA(and_value, idx, curr_group, smallest_value)

Functions

template void update_cuda< long, float > (watershed_gpu_state_t< long, float > &gpu_state, long *final_segments_d)
template void update_cuda< long, unsigned short > (watershed_gpu_state_t< long, unsigned short > &gpu_state, long *final_segments_d)
template void update_cuda< long, unsigned char > (watershed_gpu_state_t< long, unsigned char > &gpu_state, long *final_segments_d)
template void update_cuda< int, float > (watershed_gpu_state_t< int, float > &gpu_state, int *final_segments_d)
template void update_cuda< int, unsigned short > (watershed_gpu_state_t< int, unsigned short > &gpu_state, int *final_segments_d)
template void update_cuda< int, unsigned char > (watershed_gpu_state_t< int, unsigned char > &gpu_state, int *final_segments_d)
template void update_cuda< short, float > (watershed_gpu_state_t< short, float > &gpu_state, short *final_segments_d)
template void update_cuda< short, unsigned short > (watershed_gpu_state_t< short, unsigned short > &gpu_state, short *final_segments_d)
template void update_cuda< short, unsigned char > (watershed_gpu_state_t< short, unsigned char > &gpu_state, short *final_segments_d)
template<typename GROUP_T, typename IMAGE_T> __global__ void update_kernel (GROUP_T *__restrict__ current_group, GROUP_T *__restrict__ next_group, IMAGE_T *__restrict__ current_value, IMAGE_T *__restrict__ next_value, int *__restrict__ eq_and_lower, unsigned char *__restrict__ current_update, unsigned char *__restrict__ next_update, const int width, const int height, const int depth)
template<typename GROUP_T, typename IMAGE_T> void init_neighbor_offsets (watershed_gpu_state_t< GROUP_T, IMAGE_T > &gpu_state)
template<typename GROUP_T, typename IMAGE_T> __global__ void calc_neighbors_kernel (const IMAGE_T *__restrict__ image, GROUP_T *__restrict__ next_group, IMAGE_T *__restrict__ curr_value, IMAGE_T *__restrict__ next_value, int *__restrict__ eq_and_lower, const int width, const int height, const int depth)
template<typename GROUP_T, typename IMAGE_T> void set_gpu_state (state_t< GROUP_T, IMAGE_T > &state, int *eq_and_lower, watershed_gpu_state_t< GROUP_T, IMAGE_T > &gpu_state, unsigned long nVoxels)
template<typename GROUP_T, typename IMAGE_T> bool init_gpu_on_device (watershed_gpu_state_t< GROUP_T, IMAGE_T > &gpu_state, IMAGE_T *image, int imageongpu, unsigned int w, unsigned int h, unsigned int d)
template<typename GROUP_T, typename IMAGE_T> bool init_gpu (state_t< GROUP_T, IMAGE_T > &state, int *eq_and_lower, watershed_gpu_state_t< GROUP_T, IMAGE_T > &gpu_state, unsigned int w, unsigned int h, unsigned int d)
template<typename GROUP_T, typename IMAGE_T> void destroy_gpu (watershed_gpu_state_t< GROUP_T, IMAGE_T > &gpu_state)
template<typename GROUP_T, typename IMAGE_T> void update_cuda (watershed_gpu_state_t< GROUP_T, IMAGE_T > &gpu_state, GROUP_T *final_segments_d)

Variables

__constant__ int neighbor_offsets_d [19]
__device__ unsigned int changes_d

Detailed Description

CUDA-accelerated Watershed image segmentation.

Definition in file CUDAWatershed.cu.

Define Documentation

#define BLOCK_X_DIM   64
 

Definition at line 61 of file CUDAWatershed.cu.

Referenced by calc_neighbors_kernel, init_gpu, init_gpu_on_device, update_cuda, and update_kernel.

#define BLOCK_Y_DIM   2
 

Definition at line 62 of file CUDAWatershed.cu.

Referenced by calc_neighbors_kernel, init_gpu, init_gpu_on_device, update_cuda, and update_kernel.

#define BLOCK_Z_DIM   2
 

Definition at line 63 of file CUDAWatershed.cu.

Referenced by calc_neighbors_kernel, init_gpu, init_gpu_on_device, update_cuda, and update_kernel.

#define CALCULATE_NEIGHBORS_CUDA (  offset_str    ) 
 

Value:

{\
  const long idx_offset = idx + offset_str##_offset_d;\
  slope = curr_intensity - image[idx_offset];\
  if (slope < smallest_slope) {\
    smallest_slope = slope;\
    curr_lower = offset_str##_idx;\
  } else if (slope >= -FLOAT_DIFF && slope <= FLOAT_DIFF) {\
    curr_n_eq |= offset_str;\
    if (idx_offset < min_group) {\
      min_group = idx_offset;\
    }\
  } \
  }

Definition at line 205 of file CUDAWatershed.cu.

Referenced by calc_neighbors_kernel.

#define DIV_BY_32   5
 

Definition at line 157 of file CUDAWatershed.cu.

Referenced by update_kernel.

#define GIGABYTE   (1024.0f*1024.0f*1024.0f)
 

Definition at line 158 of file CUDAWatershed.cu.

#define GPU_WARP_UPDATE
 

Definition at line 50 of file CUDAWatershed.cu.

#define INST_DESTROY_GPU (  G_T    ) 
 

Value:

template void destroy_gpu<G_T,float>(watershed_gpu_state_t<G_T,float> &gpu_state);\
template void destroy_gpu<G_T,unsigned short>(watershed_gpu_state_t<G_T,unsigned short> &gpu_state);\
template void destroy_gpu<G_T,unsigned char>(watershed_gpu_state_t<G_T,unsigned char> &gpu_state);

Definition at line 74 of file CUDAWatershed.cu.

#define INST_INIT_GPU (  G_T    ) 
 

Value:

template bool init_gpu<G_T, float>(state_t<G_T,float> &state, int *eq_and_lower,\
                             watershed_gpu_state_t<G_T,float> &gpu_state,\
                             unsigned int w, unsigned int h, unsigned int d);\
template bool init_gpu<G_T, unsigned short>(state_t<G_T,unsigned short> &state, int *eq_and_lower,\
                             watershed_gpu_state_t<G_T,unsigned short> &gpu_state,\
                             unsigned int w, unsigned int h, unsigned int d);\
template bool init_gpu<G_T, unsigned char>(state_t<G_T,unsigned char> &state, int *eq_and_lower,\
                             watershed_gpu_state_t<G_T,unsigned char> &gpu_state,\
                             unsigned int w, unsigned int h, unsigned int d);

Definition at line 87 of file CUDAWatershed.cu.

#define INST_INIT_GPU_ON_DEV (  G_T    ) 
 

Value:

template bool init_gpu_on_device<G_T, float>(watershed_gpu_state_t<G_T,float> &gpu_state, float* image, int imageongpu,\
                                             unsigned int w, unsigned int h, unsigned int d);\
template bool init_gpu_on_device<G_T, unsigned short>(watershed_gpu_state_t<G_T,unsigned short> &gpu_state,\
                                unsigned short* image, int imageongpu, unsigned int w, unsigned int h, unsigned int d);\
template bool init_gpu_on_device<G_T, unsigned char>(watershed_gpu_state_t<G_T,unsigned char> &gpu_state,\
                                unsigned char* image, int imageongpu, unsigned int w, unsigned int h, unsigned int d);\

Definition at line 79 of file CUDAWatershed.cu.

#define INST_UPDATE_CUDA (  G_T    ) 
 

Value:

template void update_cuda<G_T,float>(watershed_gpu_state_t<G_T,float> &gpu_state, \
                                G_T *final_segments_d);\
template void update_cuda<G_T,unsigned short>(watershed_gpu_state_t<G_T,unsigned short> &gpu_state, \
                                G_T *final_segments_d);\
template void update_cuda<G_T,unsigned char>(watershed_gpu_state_t<G_T,unsigned char> &gpu_state, \
                                G_T *final_segments_d);

Definition at line 98 of file CUDAWatershed.cu.

#define nx_ny_offset_d   (-1 - width)
 

Definition at line 34 of file CUDAWatershed.cu.

#define nx_nz_offset_d   (-1 - heightWidth)
 

Definition at line 40 of file CUDAWatershed.cu.

#define nx_offset_d   (-1)
 

Definition at line 30 of file CUDAWatershed.cu.

#define nx_py_offset_d   (-1 + width)
 

Definition at line 35 of file CUDAWatershed.cu.

#define nx_pz_offset_d   (-1 + heightWidth)
 

Definition at line 41 of file CUDAWatershed.cu.

#define ny_nz_offset_d   (-width - heightWidth)
 

Definition at line 45 of file CUDAWatershed.cu.

#define ny_offset_d   (-width)
 

Definition at line 31 of file CUDAWatershed.cu.

#define ny_pz_offset_d   (-width + heightWidth)
 

Definition at line 46 of file CUDAWatershed.cu.

#define nz_offset_d   (-heightWidth)
 

Definition at line 32 of file CUDAWatershed.cu.

#define px_ny_offset_d   (1 - width)
 

Definition at line 37 of file CUDAWatershed.cu.

#define px_nz_offset_d   (1 - heightWidth)
 

Definition at line 42 of file CUDAWatershed.cu.

#define px_offset_d   (1)
 

Definition at line 27 of file CUDAWatershed.cu.

#define px_py_offset_d   (1 + width)
 

Definition at line 36 of file CUDAWatershed.cu.

#define px_pz_offset_d   (1 + heightWidth)
 

Definition at line 39 of file CUDAWatershed.cu.

#define py_nz_offset_d   (width - heightWidth)
 

Definition at line 47 of file CUDAWatershed.cu.

#define py_offset_d   (width)
 

Definition at line 28 of file CUDAWatershed.cu.

#define py_pz_offset_d   (width + heightWidth)
 

Definition at line 44 of file CUDAWatershed.cu.

#define pz_offset_d   (heightWidth)
 

Definition at line 29 of file CUDAWatershed.cu.

#define UPDATE_VOXEL_CUDA (  and_value,
idx,
curr_group,
smallest_value    ) 
 

Value:

{\
  if (n_eq & and_value) {\
    const long offset_idx = idx + and_value##_offset_d;\
    if (current_value[offset_idx] + FLOAT_DIFF < smallest_value) {\
      smallest_value = current_value[offset_idx];\
      offset_number = and_value##_idx;\
      next_idx = idx + and_value##_offset_d;\
    }\
    curr_group = current_group[offset_idx] < curr_group ?\
    current_group[offset_idx] : curr_group;\
  }}

Definition at line 622 of file CUDAWatershed.cu.

Referenced by update_kernel.

#define WATERSHED_INTERNAL   1
 

Definition at line 23 of file CUDAWatershed.cu.

Function Documentation

template<typename GROUP_T, typename IMAGE_T>
__global__ void calc_neighbors_kernel (  const IMAGE_T *__restrict__    image,
GROUP_T *__restrict__    next_group,
IMAGE_T *__restrict__    curr_value,
IMAGE_T *__restrict__    next_value,
int *__restrict__    eq_and_lower,
const int    width,
const int    height,
const int    depth
) 
 

Definition at line 222 of file CUDAWatershed.cu.

References BLOCK_X_DIM, BLOCK_Y_DIM, BLOCK_Z_DIM, CALCULATE_NEIGHBORS_CUDA, neighbor_offsets_d, and z.

template<typename GROUP_T, typename IMAGE_T>
void destroy_gpu (  watershed_gpu_state_t< GROUP_T, IMAGE_T > &    gpu_state ) 
 

Definition at line 500 of file CUDAWatershed.cu.

References watershed_gpu_state_t::current_update_d, watershed_gpu_state_t::current_value_d, watershed_gpu_state_t::eq_and_lower_d, watershed_gpu_state_t::init, watershed_gpu_state_t::next_update_d, watershed_gpu_state_t::next_value_d, NULL, and watershed_gpu_state_t::segments_d.

template<typename GROUP_T, typename IMAGE_T>
bool init_gpu (  state_t< GROUP_T, IMAGE_T > &    state,
int *    eq_and_lower,
watershed_gpu_state_t< GROUP_T, IMAGE_T > &    gpu_state,
unsigned int    w,
unsigned int    h,
unsigned int    d
) 
 

Definition at line 442 of file CUDAWatershed.cu.

References BLOCK_X_DIM, BLOCK_Y_DIM, BLOCK_Z_DIM, changes_d, watershed_gpu_state_t::current_update_d, watershed_gpu_state_t::current_value_d, watershed_gpu_state_t::depth, destroy_gpu, watershed_gpu_state_t::eq_and_lower_d, watershed_gpu_state_t::height, watershed_gpu_state_t::init, init_neighbor_offsets, watershed_gpu_state_t::next_update_d, watershed_gpu_state_t::next_value_d, PROFILE_MARK, watershed_gpu_state_t::segments_d, set_gpu_state, state, and watershed_gpu_state_t::width.

template<typename GROUP_T, typename IMAGE_T>
bool init_gpu_on_device (  watershed_gpu_state_t< GROUP_T, IMAGE_T > &    gpu_state,
IMAGE_T *    image,
int    imageongpu,
unsigned int    w,
unsigned int    h,
unsigned int    d
) 
 

Definition at line 362 of file CUDAWatershed.cu.

References BLOCK_X_DIM, BLOCK_Y_DIM, BLOCK_Z_DIM, changes_d, watershed_gpu_state_t::current_update_d, watershed_gpu_state_t::current_value_d, watershed_gpu_state_t::depth, destroy_gpu, watershed_gpu_state_t::eq_and_lower_d, watershed_gpu_state_t::height, watershed_gpu_state_t::init, init_neighbor_offsets, watershed_gpu_state_t::next_update_d, watershed_gpu_state_t::next_value_d, PROFILE_POP_RANGE, PROFILE_PUSH_RANGE, watershed_gpu_state_t::segments_d, and watershed_gpu_state_t::width.

template<typename GROUP_T, typename IMAGE_T>
void init_neighbor_offsets (  watershed_gpu_state_t< GROUP_T, IMAGE_T > &    gpu_state ) 
 

Definition at line 173 of file CUDAWatershed.cu.

References watershed_gpu_state_t::height, neighbor_offsets_d, nx_ny_offset, nx_nz_offset, nx_offset, nx_py_offset, nx_pz_offset, ny_nz_offset, ny_offset, ny_pz_offset, nz_offset, px_ny_offset, px_nz_offset, px_offset, px_py_offset, px_pz_offset, py_nz_offset, py_offset, py_pz_offset, pz_offset, and watershed_gpu_state_t::width.

Referenced by init_gpu, and init_gpu_on_device.

template<typename GROUP_T, typename IMAGE_T>
void set_gpu_state (  state_t< GROUP_T, IMAGE_T > &    state,
int *    eq_and_lower,
watershed_gpu_state_t< GROUP_T, IMAGE_T > &    gpu_state,
unsigned long    nVoxels
) 
 

Definition at line 346 of file CUDAWatershed.cu.

References watershed_gpu_state_t::current_value_d, watershed_gpu_state_t::eq_and_lower_d, state_t::group, watershed_gpu_state_t::next_value_d, watershed_gpu_state_t::segments_d, state, and state_t::value.

Referenced by init_gpu.

template<typename GROUP_T, typename IMAGE_T>
void update_cuda (  watershed_gpu_state_t< GROUP_T, IMAGE_T > &    gpu_state,
GROUP_T *    final_segments_d
) 
 

Definition at line 562 of file CUDAWatershed.cu.

References BLOCK_X_DIM, BLOCK_Y_DIM, BLOCK_Z_DIM, changes_d, watershed_gpu_state_t::current_update_d, watershed_gpu_state_t::current_value_d, watershed_gpu_state_t::depth, destroy_gpu, watershed_gpu_state_t::eq_and_lower_d, watershed_gpu_state_t::height, watershed_gpu_state_t::next_update_d, watershed_gpu_state_t::next_value_d, watershed_gpu_state_t::segments_d, and watershed_gpu_state_t::width.

template void update_cuda< int, float > (  watershed_gpu_state_t< int, float > &    gpu_state,
int *    final_segments_d
) 
 

template void update_cuda< int, unsigned char > (  watershed_gpu_state_t< int, unsigned char > &    gpu_state,
int *    final_segments_d
) 
 

template void update_cuda< int, unsigned short > (  watershed_gpu_state_t< int, unsigned short > &    gpu_state,
int *    final_segments_d
) 
 

template void update_cuda< long, float > (  watershed_gpu_state_t< long, float > &    gpu_state,
long *    final_segments_d
) 
 

template void update_cuda< long, unsigned char > (  watershed_gpu_state_t< long, unsigned char > &    gpu_state,
long *    final_segments_d
) 
 

template void update_cuda< long, unsigned short > (  watershed_gpu_state_t< long, unsigned short > &    gpu_state,
long *    final_segments_d
) 
 

template void update_cuda< short, float > (  watershed_gpu_state_t< short, float > &    gpu_state,
short *    final_segments_d
) 
 

template void update_cuda< short, unsigned char > (  watershed_gpu_state_t< short, unsigned char > &    gpu_state,
short *    final_segments_d
) 
 

template void update_cuda< short, unsigned short > (  watershed_gpu_state_t< short, unsigned short > &    gpu_state,
short *    final_segments_d
) 
 

template<typename GROUP_T, typename IMAGE_T>
__global__ void update_kernel (  GROUP_T *__restrict__    current_group,
GROUP_T *__restrict__    next_group,
IMAGE_T *__restrict__    current_value,
IMAGE_T *__restrict__    next_value,
int *__restrict__    eq_and_lower,
unsigned char *__restrict__    current_update,
unsigned char *__restrict__    next_update,
const int    width,
const int    height,
const int    depth
) 
 

Definition at line 636 of file CUDAWatershed.cu.

References BLOCK_X_DIM, BLOCK_Y_DIM, BLOCK_Z_DIM, changes_d, DIV_BY_32, neighbor_offsets_d, and UPDATE_VOXEL_CUDA.

Variable Documentation

__device__ unsigned int changes_d
 

Definition at line 72 of file CUDAWatershed.cu.

Referenced by init_gpu, init_gpu_on_device, update_cuda, and update_kernel.

__constant__ int neighbor_offsets_d[19]
 

Definition at line 71 of file CUDAWatershed.cu.

Referenced by calc_neighbors_kernel, init_neighbor_offsets, and update_kernel.

Generated on Thu Apr 18 02:46:02 2024 for VMD (current) by doxygen1.2.14 written by Dimitri van Heesch, © 1997-2002