#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 |
Definition in file CUDAWatershed.cu.
|
Definition at line 61 of file CUDAWatershed.cu. Referenced by calc_neighbors_kernel, init_gpu, init_gpu_on_device, update_cuda, and update_kernel. |
|
Definition at line 62 of file CUDAWatershed.cu. Referenced by calc_neighbors_kernel, init_gpu, init_gpu_on_device, update_cuda, and update_kernel. |
|
Definition at line 63 of file CUDAWatershed.cu. Referenced by calc_neighbors_kernel, init_gpu, init_gpu_on_device, update_cuda, and update_kernel. |
|
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. |
|
Definition at line 157 of file CUDAWatershed.cu. Referenced by update_kernel. |
|
Definition at line 158 of file CUDAWatershed.cu. |
|
Definition at line 50 of file CUDAWatershed.cu. |
|
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. |
|
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. |
|
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. |
|
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. |
|
Definition at line 34 of file CUDAWatershed.cu. |
|
Definition at line 40 of file CUDAWatershed.cu. |
|
Definition at line 30 of file CUDAWatershed.cu. |
|
Definition at line 35 of file CUDAWatershed.cu. |
|
Definition at line 41 of file CUDAWatershed.cu. |
|
Definition at line 45 of file CUDAWatershed.cu. |
|
Definition at line 31 of file CUDAWatershed.cu. |
|
Definition at line 46 of file CUDAWatershed.cu. |
|
Definition at line 32 of file CUDAWatershed.cu. |
|
Definition at line 37 of file CUDAWatershed.cu. |
|
Definition at line 42 of file CUDAWatershed.cu. |
|
Definition at line 27 of file CUDAWatershed.cu. |
|
Definition at line 36 of file CUDAWatershed.cu. |
|
Definition at line 39 of file CUDAWatershed.cu. |
|
Definition at line 47 of file CUDAWatershed.cu. |
|
Definition at line 28 of file CUDAWatershed.cu. |
|
Definition at line 44 of file CUDAWatershed.cu. |
|
Definition at line 29 of file CUDAWatershed.cu. |
|
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. |
|
Definition at line 23 of file CUDAWatershed.cu. |
|
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. |
|
|
|
|
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. |
|
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. |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
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. |
|
Definition at line 72 of file CUDAWatershed.cu. Referenced by init_gpu, init_gpu_on_device, update_cuda, and update_kernel. |
|
Definition at line 71 of file CUDAWatershed.cu. Referenced by calc_neighbors_kernel, init_neighbor_offsets, and update_kernel. |