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

CUDASegmentation.cu File Reference

CUDA scale-space variant of Watershed image segmentation intended for use on 3-D cryo-EM density maps. More...

#include <stdio.h>
#include "CUDAParPrefixOps.h"
#include "CUDASegmentation.h"

Go to the source code of this file.

Defines

#define CUERR
#define GROUP_BLOCK_SIZE   1024
#define VOXEL_BLOCK_SIZE   1024
#define NUM_N   18
#define COPY_CONV_INST(T)
#define INST_SEQ_GROUPS_CUDA(G_T)
#define INST_FIND_MAX_IDX_CUDA(G_T)
#define INST_HILL_CLIMB_MERGE_CUDA(G_T)
#define INST_WATERSHED_HILL_CLIMB_MERGE_CUDA(G_T)

Functions

void free_gpu_temp_storage (gpuseg_temp_storage *tmp)
template<typename GROUP_T> __global__ void update_groups_from_map (GROUP_T *__restrict__ groups_d, const GROUP_T *__restrict__ group_map_d, const long nVoxels)
template<typename GROUP_T> __global__ void init_group_map (const GROUP_T *__restrict__ groups_d, GROUP_T *__restrict__ group_map_d, const long nVoxels)
template<typename GROUP_T> __global__ void update_group_map (GROUP_T *__restrict__ group_map_d, const GROUP_T *__restrict__ scan_results, const long nGroups)
template<typename GROUP_T> long sequentialize_groups_cuda (GROUP_T *groups_d, GROUP_T *group_map_d, unsigned long nVoxels, unsigned long nGroups, gpuseg_temp_storage *tmp, gpuseg_temp_storage *scanwork)
template<typename IN_T> __device__ void getOrderedInt (IN_T input, int *retAddr)
template<> __device__ void getOrderedInt< float > (float input, int *retAddr)
template<typename IN_T> int getOrderedIntHost (IN_T floatVal)
template<> int getOrderedIntHost< float > (float floatVal)
template<typename IMAGE_T> __global__ void init_max_array (int *max_values, long nGroups)
template<typename GROUP_T, typename IMAGE_T> __global__ void find_max_values (const GROUP_T *__restrict__ groups_d, const IMAGE_T *__restrict__ image_d, int *__restrict__ max_values, const long nVoxels)
template<typename GROUP_T, typename IMAGE_T> __global__ void find_max_values_warpred (const GROUP_T *__restrict__ groups_d, const IMAGE_T *__restrict__ image_d, int *__restrict__ max_values, const long nVoxels)
template<typename GROUP_T, typename IMAGE_T> __global__ void find_max_values_shm (const GROUP_T *__restrict__ groups_d, const IMAGE_T *__restrict__ image_d, int *__restrict__ max_values, const long nVoxels, int initval, long nGroups)
template<typename GROUP_T, typename IMAGE_T> __global__ void find_max_values_shm_2xl (const GROUP_T *__restrict__ groups_d, const IMAGE_T *__restrict__ image_d, int *__restrict__ max_values, const long nVoxels, int initval, long nGroups)
template<typename GROUP_T, typename IMAGE_T> __global__ void find_max_values_shm_4xl (const GROUP_T *__restrict__ groups_d, const IMAGE_T *__restrict__ image_d, int *__restrict__ max_values, const long nVoxels, int initval, long nGroups)
template<typename GROUP_T, typename IMAGE_T> __global__ void update_group_idx (const GROUP_T *__restrict__ groups_d, const IMAGE_T *__restrict__ image_d, const int *__restrict max_values, unsigned long *__restrict__ group_idx, const unsigned long nVoxels)
template<typename GROUP_T, typename IMAGE_T> void find_groups_max_idx_cuda (GROUP_T *groups_d, IMAGE_T *image_d, unsigned long *max_idx, unsigned long nVoxels, unsigned long nGroups, gpuseg_temp_storage *tmp)
template<typename GROUP_T, typename IMAGE_T> __global__ void hill_climb_kernel (IMAGE_T *image_d, unsigned long *max_int_d, GROUP_T *group_map_d, GROUP_T *groups_d, int height, int width, int depth, unsigned long nGroups)
template<typename GROUP_T, typename IMAGE_T> void hill_climb_merge_cuda (GROUP_T *groups_d, IMAGE_T *image_d, unsigned long *max_idx_d, GROUP_T *group_map_d, int height, int width, int depth, unsigned long nGroups)
template<typename GROUP_T, typename IMAGE_T> void watershed_hill_climb_merge_cuda (GROUP_T *segments_d, GROUP_T *new_segments_d, IMAGE_T *image_d, GROUP_T *group_map_d, unsigned long *max_idx_d, long height, long width, long depth, unsigned long nGroups)
template<typename IN_T, typename OUT_T> __global__ void copy_and_convert_kernel (IN_T *in, OUT_T *out, long num_elements)
template<typename IN_T, typename OUT_T> void copy_and_convert_type_cuda (IN_T *in, OUT_T *out, long num_elements)

Variables

__device__ unsigned long nGroups_d

Detailed Description

CUDA scale-space variant of Watershed image segmentation intended for use on 3-D cryo-EM density maps.

Definition in file CUDASegmentation.cu.

Define Documentation

#define COPY_CONV_INST (  T    ) 
 

Value:

template void copy_and_convert_type_cuda<T,unsigned long>(T* in, unsigned long* out, long num_elements);\
  template void copy_and_convert_type_cuda<T,unsigned int>(T* in, unsigned int* out, long num_elements);\
  template void copy_and_convert_type_cuda<T,unsigned short>(T* in, unsigned short* out, long num_elements);\
  template void copy_and_convert_type_cuda<T,unsigned char>(T* in, unsigned char* out, long num_elements);\
  template void copy_and_convert_type_cuda<T,long>(T* in, long* out, long num_elements);\
  template void copy_and_convert_type_cuda<T,int>(T* in, int* out, long num_elements);\
  template void copy_and_convert_type_cuda<T,short>(T* in, short* out, long num_elements);\
  template void copy_and_convert_type_cuda<T,char>(T* in, char* out, long num_elements);\
  template void copy_and_convert_type_cuda<T,float>(T* in, float* out, long num_elements);\

Definition at line 631 of file CUDASegmentation.cu.

#define CUERR
 

Definition at line 33 of file CUDASegmentation.cu.

Referenced by find_groups_max_idx_cuda, hill_climb_merge_cuda, sequentialize_groups_cuda, and watershed_hill_climb_merge_cuda.

#define GROUP_BLOCK_SIZE   1024
 

Definition at line 36 of file CUDASegmentation.cu.

Referenced by find_groups_max_idx_cuda, and sequentialize_groups_cuda.

#define INST_FIND_MAX_IDX_CUDA (  G_T    ) 
 

Value:

template void find_groups_max_idx_cuda<G_T, float>(G_T* groups_d, float* image_d, unsigned long* max_idx,\
                                                   unsigned long nVoxels, unsigned long nGroups,\
                                                   gpuseg_temp_storage *tmp);\
template void find_groups_max_idx_cuda<G_T, unsigned short>(G_T* groups_d, unsigned short* image_d, unsigned long* max_idx,\
                                                   unsigned long nVoxels, unsigned long nGroups,\
                                                   gpuseg_temp_storage *tmp);\
template void find_groups_max_idx_cuda<G_T, unsigned char>(G_T* groups_d, unsigned char* image_d, unsigned long* max_idx,\
                                                   unsigned long nVoxels, unsigned long nGroups,\
                                                   gpuseg_temp_storage *tmp);

Definition at line 648 of file CUDASegmentation.cu.

#define INST_HILL_CLIMB_MERGE_CUDA (  G_T    ) 
 

Value:

template void hill_climb_merge_cuda<G_T, float>(G_T* groups_d, float* image_d, unsigned long* max_idx_d,\
                          G_T* group_map_d, int height, int width, int depth, unsigned long nGroups);\
template void hill_climb_merge_cuda<G_T, unsigned short>(G_T* groups_d, unsigned short* image_d, unsigned long* max_idx_d,\
                          G_T* group_map_d, int height, int width, int depth, unsigned long nGroups);\
template void hill_climb_merge_cuda<G_T, unsigned char>(G_T* groups_d, unsigned char* image_d, unsigned long* max_idx_d,\
                          G_T* group_map_d, int height, int width, int depth, unsigned long nGroups);

Definition at line 659 of file CUDASegmentation.cu.

#define INST_SEQ_GROUPS_CUDA (  G_T    ) 
 

Value:

template long sequentialize_groups_cuda<G_T>(G_T* groups_d, G_T* group_map_d,\
                                         unsigned long nVoxels, unsigned long nGroups,\
                                         gpuseg_temp_storage *tmp,\
                                         gpuseg_temp_storage *scanwork);\

Definition at line 642 of file CUDASegmentation.cu.

#define INST_WATERSHED_HILL_CLIMB_MERGE_CUDA (  G_T    ) 
 

Value:

template void watershed_hill_climb_merge_cuda<G_T, float>(G_T* segments_d, G_T* new_segments_d, float* image_d, G_T* group_map_d,\
                                     unsigned long* max_idx_d, long height, long width, long depth, unsigned long nGroups);\
template void watershed_hill_climb_merge_cuda<G_T, unsigned short>(G_T* segments_d, G_T* new_segments_d,\
                                     unsigned short* image_d, G_T* group_map_d,\
                                     unsigned long* max_idx_d, long height, long width, long depth, unsigned long nGroups);\
template void watershed_hill_climb_merge_cuda<G_T, unsigned char>(G_T* segments_d, G_T* new_segments_d,\
                                     unsigned char* image_d, G_T* group_map_d,\
                                     unsigned long* max_idx_d, long height, long width, long depth, unsigned long nGroups);

Definition at line 667 of file CUDASegmentation.cu.

#define NUM_N   18
 

Definition at line 486 of file CUDASegmentation.cu.

Referenced by hill_climb_kernel.

#define VOXEL_BLOCK_SIZE   1024
 

Definition at line 37 of file CUDASegmentation.cu.

Referenced by copy_and_convert_type_cuda, find_groups_max_idx_cuda, hill_climb_merge_cuda, sequentialize_groups_cuda, and watershed_hill_climb_merge_cuda.

Function Documentation

template<typename IN_T, typename OUT_T>
__global__ void copy_and_convert_kernel (  IN_T *    in,
OUT_T *    out,
long    num_elements
) 
 

Definition at line 616 of file CUDASegmentation.cu.

template<typename IN_T, typename OUT_T>
void copy_and_convert_type_cuda (  IN_T *    in,
OUT_T *    out,
long    num_elements
) 
 

Definition at line 624 of file CUDASegmentation.cu.

References VOXEL_BLOCK_SIZE.

template<typename GROUP_T, typename IMAGE_T>
void find_groups_max_idx_cuda (  GROUP_T *    groups_d,
IMAGE_T *    image_d,
unsigned long *    max_idx,
unsigned long    nVoxels,
unsigned long    nGroups,
gpuseg_temp_storage *    tmp
) 
 

Definition at line 420 of file CUDASegmentation.cu.

References CUERR, getOrderedIntHost, GROUP_BLOCK_SIZE, NULL, gpuseg_temp_storage::sz, gpuseg_temp_storage::tmp_d, and VOXEL_BLOCK_SIZE.

template<typename GROUP_T, typename IMAGE_T>
__global__ void find_max_values (  const GROUP_T *__restrict__    groups_d,
const IMAGE_T *__restrict__    image_d,
int *__restrict__    max_values,
const long    nVoxels
) 
 

Definition at line 226 of file CUDASegmentation.cu.

References getOrderedInt.

template<typename GROUP_T, typename IMAGE_T>
__global__ void find_max_values_shm (  const GROUP_T *__restrict__    groups_d,
const IMAGE_T *__restrict__    image_d,
int *__restrict__    max_values,
const long    nVoxels,
int    initval,
long    nGroups
) 
 

Definition at line 288 of file CUDASegmentation.cu.

References getOrderedInt.

template<typename GROUP_T, typename IMAGE_T>
__global__ void find_max_values_shm_2xl (  const GROUP_T *__restrict__    groups_d,
const IMAGE_T *__restrict__    image_d,
int *__restrict__    max_values,
const long    nVoxels,
int    initval,
long    nGroups
) 
 

Definition at line 317 of file CUDASegmentation.cu.

References getOrderedInt.

template<typename GROUP_T, typename IMAGE_T>
__global__ void find_max_values_shm_4xl (  const GROUP_T *__restrict__    groups_d,
const IMAGE_T *__restrict__    image_d,
int *__restrict__    max_values,
const long    nVoxels,
int    initval,
long    nGroups
) 
 

Definition at line 353 of file CUDASegmentation.cu.

References getOrderedInt.

template<typename GROUP_T, typename IMAGE_T>
__global__ void find_max_values_warpred (  const GROUP_T *__restrict__    groups_d,
const IMAGE_T *__restrict__    image_d,
int *__restrict__    max_values,
const long    nVoxels
) 
 

Definition at line 244 of file CUDASegmentation.cu.

References getOrderedInt.

void free_gpu_temp_storage (  gpuseg_temp_storage *    tmp ) 
 

Definition at line 45 of file CUDASegmentation.cu.

References NULL, gpuseg_temp_storage::sz, and gpuseg_temp_storage::tmp_d.

template<typename IN_T>
__device__ void getOrderedInt (  IN_T    input,
int *    retAddr
) 
 

Definition at line 188 of file CUDASegmentation.cu.

Referenced by find_max_values, find_max_values_shm, find_max_values_shm_2xl, find_max_values_shm_4xl, find_max_values_warpred, init_max_array, and update_group_idx.

template<>
__device__ void getOrderedInt< float > (  float    input,
int *    retAddr
) 
 

template<typename IN_T>
int getOrderedIntHost (  IN_T    floatVal ) 
 

Definition at line 202 of file CUDASegmentation.cu.

Referenced by find_groups_max_idx_cuda.

template<>
int getOrderedIntHost< float > (  float    floatVal ) 
 

template<typename GROUP_T, typename IMAGE_T>
__global__ void hill_climb_kernel (  IMAGE_T *    image_d,
unsigned long *    max_int_d,
GROUP_T *    group_map_d,
GROUP_T *    groups_d,
int    height,
int    width,
int    depth,
unsigned long    nGroups
) 
 

Definition at line 489 of file CUDASegmentation.cu.

References NUM_N, and z.

template<typename GROUP_T, typename IMAGE_T>
void hill_climb_merge_cuda (  GROUP_T *    groups_d,
IMAGE_T *    image_d,
unsigned long *    max_idx_d,
GROUP_T *    group_map_d,
int    height,
int    width,
int    depth,
unsigned long    nGroups
) 
 

Definition at line 580 of file CUDASegmentation.cu.

References CUERR, and VOXEL_BLOCK_SIZE.

template<typename GROUP_T>
__global__ void init_group_map (  const GROUP_T *__restrict__    groups_d,
GROUP_T *__restrict__    group_map_d,
const long    nVoxels
) 
 

Definition at line 67 of file CUDASegmentation.cu.

template<typename IMAGE_T>
__global__ void init_max_array (  int *    max_values,
long    nGroups
) 
 

Definition at line 216 of file CUDASegmentation.cu.

References getOrderedInt.

template<typename GROUP_T>
long sequentialize_groups_cuda (  GROUP_T *    groups_d,
GROUP_T *    group_map_d,
unsigned long    nVoxels,
unsigned long    nGroups,
gpuseg_temp_storage *    tmp,
gpuseg_temp_storage *    scanwork
) 
 

Definition at line 96 of file CUDASegmentation.cu.

References CUERR, dev_excl_scan_sum, dev_excl_scan_sum_tmpsz, GROUP_BLOCK_SIZE, nGroups_d, NULL, gpuseg_temp_storage::sz, gpuseg_temp_storage::tmp_d, and VOXEL_BLOCK_SIZE.

template<typename GROUP_T, typename IMAGE_T>
__global__ void update_group_idx (  const GROUP_T *__restrict__    groups_d,
const IMAGE_T *__restrict__    image_d,
const int *__restrict    max_values,
unsigned long *__restrict__    group_idx,
const unsigned long    nVoxels
) 
 

Definition at line 402 of file CUDASegmentation.cu.

References getOrderedInt.

template<typename GROUP_T>
__global__ void update_group_map (  GROUP_T *__restrict__    group_map_d,
const GROUP_T *__restrict__    scan_results,
const long    nGroups
) 
 

Definition at line 79 of file CUDASegmentation.cu.

References nGroups_d.

template<typename GROUP_T>
__global__ void update_groups_from_map (  GROUP_T *__restrict__    groups_d,
const GROUP_T *__restrict__    group_map_d,
const long    nVoxels
) 
 

Definition at line 55 of file CUDASegmentation.cu.

template<typename GROUP_T, typename IMAGE_T>
void watershed_hill_climb_merge_cuda (  GROUP_T *    segments_d,
GROUP_T *    new_segments_d,
IMAGE_T *    image_d,
GROUP_T *    group_map_d,
unsigned long *    max_idx_d,
long    height,
long    width,
long    depth,
unsigned long    nGroups
) 
 

Definition at line 598 of file CUDASegmentation.cu.

References CUERR, and VOXEL_BLOCK_SIZE.

Variable Documentation

__device__ unsigned long nGroups_d
 

Definition at line 43 of file CUDASegmentation.cu.

Referenced by sequentialize_groups_cuda, and update_group_map.

Generated on Fri Nov 8 02:45:58 2024 for VMD (current) by doxygen1.2.14 written by Dimitri van Heesch, © 1997-2002