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

CUDAGaussianBlur.cu File Reference

CUDA kernels for Gaussian blur filtering of 3-D volumes. More...

#include <stdio.h>
#include "CUDAGaussianBlur.h"
#include "Watershed.h"
#include "ProfileHooks.h"
#include <cuda_fp16.h>

Go to the source code of this file.

Defines

#define WATERSHED_INTERNAL   1
#define CUERR
#define KERNEL_MAX_SIZE   4096
#define BLOCKDIM_X   128
#define BLOCKDIM_Y   1
#define BLOCKDIM_Z   1
#define XCONV_BLOCKDIM_X   96
#define XCONV_BLOCKDIM_Y   1
#define XCONV_BLOCKDIM_Z   1
#define YCONV_BLOCKDIM_X   16
#define YCONV_BLOCKDIM_Y   8
#define YCONV_BLOCKDIM_Z   1
#define ZCONV_BLOCKDIM_X   4
#define ZCONV_BLOCKDIM_Y   1
#define ZCONV_BLOCKDIM_Z   32
#define X_PADDING   32
#define SWAPF(a, b)
#define INST_GAUSSIAN_CUDA(I_T)

Functions

template<typename T> __device__ __forceinline__
void 		convert_type_and_assign_val (T *arr, long idx, float val)
template<> __device__ __forceinline__
void 		convert_type_and_assign_val< float > (float *arr, long idx, float val)
void copy_array_from_gpu (void *arr, void *arr_d, int bytes)
void copy_array_to_gpu (void *arr_d, void *arr, int bytes)
void * alloc_cuda_array (int bytes)
void free_cuda_array (void *arr)
void set_gaussian_1D_kernel_cuda (float *kernel, int kernel_size)
template<typename IMAGE_T, const int offset> __global__ void convolution_x_kernel_s (const IMAGE_T *__restrict__ src_d, IMAGE_T *__restrict__ dst_d, int width, int height, int depth)
template<typename IMAGE_T, const int offset> __global__ void convolution_y_kernel_s (const IMAGE_T *__restrict__ src_d, IMAGE_T *__restrict__ dst_d, int width, int height, int depth)
template<typename IMAGE_T, const int offset> __global__ void convolution_z_kernel_s (const IMAGE_T *__restrict__ src_d, IMAGE_T *__restrict__ dst_d, int width, int height, int depth)
template<typename IMAGE_T, const int offset> __global__ void convolution_x_kernel (const IMAGE_T *__restrict__ src_d, IMAGE_T *__restrict__ dst_d, int width, int height, int depth)
template<typename IMAGE_T> __global__ void convolution_x_kernel (const IMAGE_T *__restrict__ src_d, IMAGE_T *__restrict__ dst_d, int offset, int width, int height, int depth)
template<typename IMAGE_T, const int offset> __global__ void convolution_y_kernel (const IMAGE_T *__restrict__ src_d, IMAGE_T *__restrict__ dst_d, int width, int height, int depth)
template<typename IMAGE_T> __global__ void convolution_y_kernel (const IMAGE_T *__restrict__ src_d, IMAGE_T *__restrict__ dst_d, int offset, int width, int height, int depth)
template<typename IMAGE_T> __global__ void convolution_z_kernel (const IMAGE_T *__restrict__ src_d, IMAGE_T *__restrict__ dst_d, int offset, int width, int height, int depth)
template<typename IMAGE_T, const int offset> __global__ void convolution_z_kernel (const IMAGE_T *__restrict__ src_d, IMAGE_T *__restrict__ dst_d, int width, int height, int depth)
template<typename IMAGE_T> void gaussian1D_x_cuda (IMAGE_T *src_d, IMAGE_T *dst_d, int kernel_size, int width, int height, int depth)
template<typename IMAGE_T> void gaussian1D_y_cuda (IMAGE_T *src_d, IMAGE_T *dst_d, int kernel_size, int width, int height, int depth)
template<typename IMAGE_T> void gaussian1D_z_cuda (IMAGE_T *src_d, IMAGE_T *dst_d, int kernel_size, int width, int height, int depth)

Variables

__constant__ float kernel_d_c [KERNEL_MAX_SIZE]

Detailed Description

CUDA kernels for Gaussian blur filtering of 3-D volumes.

Primarily for use in the context of scale-space filtering used for 3-D image segmentation of Cryo-EM density maps.

Definition in file CUDAGaussianBlur.cu.

Define Documentation

#define BLOCKDIM_X   128
 

Definition at line 46 of file CUDAGaussianBlur.cu.

#define BLOCKDIM_Y   1
 

Definition at line 47 of file CUDAGaussianBlur.cu.

#define BLOCKDIM_Z   1
 

Definition at line 48 of file CUDAGaussianBlur.cu.

#define CUERR
 

Definition at line 39 of file CUDAGaussianBlur.cu.

Referenced by alloc_cuda_array, copy_array_from_gpu, copy_array_to_gpu, gaussian1D_x_cuda, gaussian1D_y_cuda, gaussian1D_z_cuda, and set_gaussian_1D_kernel_cuda.

#define INST_GAUSSIAN_CUDA (  I_T    ) 
 

Value:

template void gaussian1D_x_cuda<I_T>(I_T* src_d, I_T* dst_d, int kernel_size,\
                                       int width, int height, int depth);\
template void gaussian1D_y_cuda<I_T>(I_T* src_d, I_T* dst_d, int kernel_size,\
                                       int width, int height, int depth);\
template void gaussian1D_z_cuda<I_T>(I_T* src_d, I_T* dst_d, int kernel_size,\
                                       int width, int height, int depth);

Definition at line 71 of file CUDAGaussianBlur.cu.

#define KERNEL_MAX_SIZE   4096
 

Definition at line 42 of file CUDAGaussianBlur.cu.

Referenced by set_gaussian_1D_kernel_cuda.

#define SWAPF (  a,
b    ) 
 

Value:

{\
  float* t = a;\
  a = b;\
  b = t;\
}

Definition at line 65 of file CUDAGaussianBlur.cu.

#define WATERSHED_INTERNAL   1
 

Definition at line 24 of file CUDAGaussianBlur.cu.

#define X_PADDING   32
 

Definition at line 63 of file CUDAGaussianBlur.cu.

#define XCONV_BLOCKDIM_X   96
 

Definition at line 51 of file CUDAGaussianBlur.cu.

Referenced by convolution_x_kernel, convolution_x_kernel_s, and gaussian1D_x_cuda.

#define XCONV_BLOCKDIM_Y   1
 

Definition at line 52 of file CUDAGaussianBlur.cu.

Referenced by convolution_x_kernel, convolution_x_kernel_s, and gaussian1D_x_cuda.

#define XCONV_BLOCKDIM_Z   1
 

Definition at line 53 of file CUDAGaussianBlur.cu.

Referenced by convolution_x_kernel, convolution_x_kernel_s, and gaussian1D_x_cuda.

#define YCONV_BLOCKDIM_X   16
 

Definition at line 55 of file CUDAGaussianBlur.cu.

Referenced by convolution_y_kernel, convolution_y_kernel_s, and gaussian1D_y_cuda.

#define YCONV_BLOCKDIM_Y   8
 

Definition at line 56 of file CUDAGaussianBlur.cu.

Referenced by convolution_y_kernel, convolution_y_kernel_s, and gaussian1D_y_cuda.

#define YCONV_BLOCKDIM_Z   1
 

Definition at line 57 of file CUDAGaussianBlur.cu.

Referenced by convolution_y_kernel, convolution_y_kernel_s, and gaussian1D_y_cuda.

#define ZCONV_BLOCKDIM_X   4
 

Definition at line 59 of file CUDAGaussianBlur.cu.

Referenced by convolution_z_kernel, convolution_z_kernel_s, and gaussian1D_z_cuda.

#define ZCONV_BLOCKDIM_Y   1
 

Definition at line 60 of file CUDAGaussianBlur.cu.

Referenced by convolution_z_kernel, convolution_z_kernel_s, and gaussian1D_z_cuda.

#define ZCONV_BLOCKDIM_Z   32
 

Definition at line 61 of file CUDAGaussianBlur.cu.

Referenced by convolution_z_kernel, convolution_z_kernel_s, and gaussian1D_z_cuda.

Function Documentation

void* alloc_cuda_array (  int    bytes ) 
 

Definition at line 118 of file CUDAGaussianBlur.cu.

References CUERR.

template<typename T>
__device__ __forceinline__ void convert_type_and_assign_val (  T *    arr,
long    idx,
float    val
) 
 

Definition at line 87 of file CUDAGaussianBlur.cu.

Referenced by convolution_x_kernel, convolution_x_kernel_s, convolution_y_kernel, convolution_y_kernel_s, convolution_z_kernel, and convolution_z_kernel_s.

template<>
__device__ __forceinline__ void convert_type_and_assign_val< float > (  float *    arr,
long    idx,
float    val
) 
 

template<typename IMAGE_T>
__global__ void convolution_x_kernel (  const IMAGE_T *__restrict__    src_d,
IMAGE_T *__restrict__    dst_d,
int    offset,
int    width,
int    height,
int    depth
) 
 

Definition at line 367 of file CUDAGaussianBlur.cu.

References convert_type_and_assign_val, kernel_d_c, XCONV_BLOCKDIM_X, XCONV_BLOCKDIM_Y, XCONV_BLOCKDIM_Z, and z.

template<typename IMAGE_T, const int offset>
__global__ void convolution_x_kernel (  const IMAGE_T *__restrict__    src_d,
IMAGE_T *__restrict__    dst_d,
int    width,
int    height,
int    depth
) 
 

Definition at line 324 of file CUDAGaussianBlur.cu.

References convert_type_and_assign_val, kernel_d_c, XCONV_BLOCKDIM_X, XCONV_BLOCKDIM_Y, XCONV_BLOCKDIM_Z, and z.

template<typename IMAGE_T, const int offset>
__global__ void convolution_x_kernel_s (  const IMAGE_T *__restrict__    src_d,
IMAGE_T *__restrict__    dst_d,
int    width,
int    height,
int    depth
) 
 

Definition at line 144 of file CUDAGaussianBlur.cu.

References convert_type_and_assign_val, kernel_d_c, XCONV_BLOCKDIM_X, XCONV_BLOCKDIM_Y, XCONV_BLOCKDIM_Z, and z.

Referenced by gaussian1D_x_cuda.

template<typename IMAGE_T>
__global__ void convolution_y_kernel (  const IMAGE_T *__restrict__    src_d,
IMAGE_T *__restrict__    dst_d,
int    offset,
int    width,
int    height,
int    depth
) 
 

Definition at line 446 of file CUDAGaussianBlur.cu.

References convert_type_and_assign_val, kernel_d_c, YCONV_BLOCKDIM_X, YCONV_BLOCKDIM_Y, YCONV_BLOCKDIM_Z, and z.

template<typename IMAGE_T, const int offset>
__global__ void convolution_y_kernel (  const IMAGE_T *__restrict__    src_d,
IMAGE_T *__restrict__    dst_d,
int    width,
int    height,
int    depth
) 
 

Definition at line 403 of file CUDAGaussianBlur.cu.

References convert_type_and_assign_val, kernel_d_c, YCONV_BLOCKDIM_X, YCONV_BLOCKDIM_Y, YCONV_BLOCKDIM_Z, and z.

template<typename IMAGE_T, const int offset>
__global__ void convolution_y_kernel_s (  const IMAGE_T *__restrict__    src_d,
IMAGE_T *__restrict__    dst_d,
int    width,
int    height,
int    depth
) 
 

Definition at line 204 of file CUDAGaussianBlur.cu.

References convert_type_and_assign_val, kernel_d_c, YCONV_BLOCKDIM_X, YCONV_BLOCKDIM_Y, YCONV_BLOCKDIM_Z, and z.

Referenced by gaussian1D_y_cuda.

template<typename IMAGE_T, const int offset>
__global__ void convolution_z_kernel (  const IMAGE_T *__restrict__    src_d,
IMAGE_T *__restrict__    dst_d,
int    width,
int    height,
int    depth
) 
 

Definition at line 519 of file CUDAGaussianBlur.cu.

References convert_type_and_assign_val, kernel_d_c, z, ZCONV_BLOCKDIM_X, ZCONV_BLOCKDIM_Y, and ZCONV_BLOCKDIM_Z.

template<typename IMAGE_T>
__global__ void convolution_z_kernel (  const IMAGE_T *__restrict__    src_d,
IMAGE_T *__restrict__    dst_d,
int    offset,
int    width,
int    height,
int    depth
) 
 

Definition at line 482 of file CUDAGaussianBlur.cu.

References convert_type_and_assign_val, kernel_d_c, z, ZCONV_BLOCKDIM_X, ZCONV_BLOCKDIM_Y, and ZCONV_BLOCKDIM_Z.

template<typename IMAGE_T, const int offset>
__global__ void convolution_z_kernel_s (  const IMAGE_T *__restrict__    src_d,
IMAGE_T *__restrict__    dst_d,
int    width,
int    height,
int    depth
) 
 

Definition at line 263 of file CUDAGaussianBlur.cu.

References convert_type_and_assign_val, kernel_d_c, z, ZCONV_BLOCKDIM_X, ZCONV_BLOCKDIM_Y, and ZCONV_BLOCKDIM_Z.

Referenced by gaussian1D_z_cuda.

void copy_array_from_gpu (  void *    arr,
void *    arr_d,
int    bytes
) 
 

Definition at line 102 of file CUDAGaussianBlur.cu.

References CUERR, PROFILE_POP_RANGE, and PROFILE_PUSH_RANGE.

void copy_array_to_gpu (  void *    arr_d,
void *    arr,
int    bytes
) 
 

Definition at line 112 of file CUDAGaussianBlur.cu.

References CUERR.

void free_cuda_array (  void *    arr ) 
 

Definition at line 126 of file CUDAGaussianBlur.cu.

template<typename IMAGE_T>
void gaussian1D_x_cuda (  IMAGE_T *    src_d,
IMAGE_T *    dst_d,
int    kernel_size,
int    width,
int    height,
int    depth
) 
 

Definition at line 564 of file CUDAGaussianBlur.cu.

References convolution_x_kernel_s, CUERR, XCONV_BLOCKDIM_X, XCONV_BLOCKDIM_Y, and XCONV_BLOCKDIM_Z.

template<typename IMAGE_T>
void gaussian1D_y_cuda (  IMAGE_T *    src_d,
IMAGE_T *    dst_d,
int    kernel_size,
int    width,
int    height,
int    depth
) 
 

Definition at line 637 of file CUDAGaussianBlur.cu.

References convolution_y_kernel_s, CUERR, YCONV_BLOCKDIM_X, YCONV_BLOCKDIM_Y, and YCONV_BLOCKDIM_Z.

template<typename IMAGE_T>
void gaussian1D_z_cuda (  IMAGE_T *    src_d,
IMAGE_T *    dst_d,
int    kernel_size,
int    width,
int    height,
int    depth
) 
 

Definition at line 697 of file CUDAGaussianBlur.cu.

References convolution_z_kernel_s, CUERR, ZCONV_BLOCKDIM_X, ZCONV_BLOCKDIM_Y, and ZCONV_BLOCKDIM_Z.

void set_gaussian_1D_kernel_cuda (  float *    kernel,
int    kernel_size
) 
 

Definition at line 131 of file CUDAGaussianBlur.cu.

References CUERR, kernel_d_c, and KERNEL_MAX_SIZE.

Variable Documentation

__constant__ float kernel_d_c[KERNEL_MAX_SIZE]
 

Definition at line 44 of file CUDAGaussianBlur.cu.

Referenced by convolution_x_kernel, convolution_x_kernel_s, convolution_y_kernel, convolution_y_kernel_s, convolution_z_kernel, convolution_z_kernel_s, and set_gaussian_1D_kernel_cuda.

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